Compare commits
5
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 SolitaryThinker
|
0a231e0a1a | ||
|
SolitaryThinker
|
e9efd52211 | ||
|
SolitaryThinker
|
2b3109b68d | ||
|
SolitaryThinker
|
ebd4d59951 | ||
|
Aryan Kumar
|
a6eb382d34 |
@@ -481,7 +481,6 @@ steps:
|
||||
- "fastvideo/layers/**"
|
||||
- "fastvideo/worker/**"
|
||||
- "fastvideo/entrypoints/**"
|
||||
- "fastvideo/performance/**"
|
||||
- "fastvideo/tests/performance/**"
|
||||
- ".buildkite/performance-benchmarks/**"
|
||||
- "pyproject.toml"
|
||||
|
||||
@@ -1,133 +0,0 @@
|
||||
#!/usr/bin/env bash
|
||||
# Gate the expensive Buildkite full suite on the cheap GitHub checks.
|
||||
#
|
||||
# Polls the workflow runs for the PR head commit and only exits 0 once the
|
||||
# watched cheap workflows (pre-commit, docs build) have succeeded, so the
|
||||
# 'ready' label cannot burn ~20 GPU lanes on a head that a cheap check has
|
||||
# already doomed.
|
||||
#
|
||||
# Semantics:
|
||||
# - watched run completed with a bad conclusion -> exit 1 (fail CLOSED:
|
||||
# no full suite; the next push re-arms via the 'synchronize' trigger)
|
||||
# - watched run cancelled -> still pending: the docs
|
||||
# workflow's repo-global 'pages' concurrency group cancels runs superseded
|
||||
# by unrelated pushes, so 'cancelled' is not a verdict on this PR
|
||||
# - watched runs pending -> poll until done
|
||||
# - docs run absent -> not applicable after a
|
||||
# short grace period ('Deploy Documentation' is path-filtered on PRs)
|
||||
# - pre-commit run absent -> keep polling: pre-commit
|
||||
# is never path-filtered, so its absence is always anomalous
|
||||
# - 'ready' label removed while waiting -> exit 1 (fail CLOSED:
|
||||
# un-labeling is a deliberate maintainer action)
|
||||
# - GitHub API unreachable or timeout -> exit 0 (fail OPEN,
|
||||
# loud warning: never brick CI on a GitHub outage)
|
||||
#
|
||||
# Required env: PR_SHA (PR head commit), PR_NUMBER, GITHUB_REPOSITORY, GH_TOKEN.
|
||||
set -euo pipefail
|
||||
|
||||
: "${PR_SHA:?PR_SHA (PR head commit) is required}"
|
||||
: "${PR_NUMBER:?PR_NUMBER (pull request number) is required}"
|
||||
: "${GITHUB_REPOSITORY:?GITHUB_REPOSITORY is required}"
|
||||
|
||||
# Workflow-level `name:` values that must be green before the full suite
|
||||
# may start. "Deploy Documentation" is path-filtered on PRs, so its run may
|
||||
# legitimately never exist; pre-commit always runs, so it must appear.
|
||||
WATCHED_NAMES='["pre-commit", "Deploy Documentation"]'
|
||||
WATCHED_REGEX='^(pre-commit|Deploy Documentation)$'
|
||||
POLL_SECS="${POLL_SECS:-20}"
|
||||
GRACE_SECS="${GRACE_SECS:-60}"
|
||||
MAX_WAIT_SECS="${MAX_WAIT_SECS:-1500}"
|
||||
|
||||
# Bound each API call so a hung connection hits the 3-strike fail-open path
|
||||
# instead of pinning the loop until the job timeout (which would fail closed
|
||||
# on exactly the GitHub-outage case this script is meant to survive).
|
||||
if command -v timeout >/dev/null 2>&1; then
|
||||
gh_api() { timeout 30 gh api "$@"; }
|
||||
else
|
||||
gh_api() { gh api "$@"; } # macOS dev boxes; CI always has coreutils timeout
|
||||
fi
|
||||
|
||||
# The workflow checked the label before starting the gate, but the wait can
|
||||
# last ~25 min: re-check once before any exit 0 and fail closed if 'ready'
|
||||
# was removed in the meantime. An API error here proceeds (the label was
|
||||
# present when the gate started; never brick CI on an outage).
|
||||
recheck_ready_label() {
|
||||
local pr_json
|
||||
if pr_json=$(gh_api "repos/${GITHUB_REPOSITORY}/pulls/${PR_NUMBER}" 2>/dev/null); then
|
||||
if ! jq -e '[.labels[]?.name] | index("ready")' <<<"$pr_json" >/dev/null 2>&1; then
|
||||
echo "::error::PR #${PR_NUMBER} no longer has the 'ready' label —" \
|
||||
"NOT triggering the Buildkite full suite. Re-add the label to re-arm."
|
||||
exit 1
|
||||
fi
|
||||
else
|
||||
echo "::warning::Could not re-check the 'ready' label on PR #${PR_NUMBER}; proceeding (it was present when the gate started)."
|
||||
fi
|
||||
}
|
||||
|
||||
start=$(date +%s)
|
||||
api_fails=0
|
||||
missing=""
|
||||
|
||||
while true; do
|
||||
elapsed=$(( $(date +%s) - start ))
|
||||
|
||||
if runs_json=$(gh_api "repos/${GITHUB_REPOSITORY}/actions/runs?head_sha=${PR_SHA}&per_page=100" 2>/dev/null) \
|
||||
&& state=$(jq --arg re "$WATCHED_REGEX" '
|
||||
[.workflow_runs[]? | select(.name // "" | test($re))]
|
||||
| group_by(.name) | map(max_by(.id))
|
||||
| map({name, status, conclusion})' <<<"$runs_json" 2>/dev/null); then
|
||||
api_fails=0
|
||||
echo "t+${elapsed}s watched checks: $(jq -c . <<<"$state")"
|
||||
|
||||
failed=$(jq -r '[.[] | select(.status == "completed"
|
||||
and (.conclusion | IN("success", "skipped", "neutral", "cancelled") | not))]
|
||||
| map(.name) | join(", ")' <<<"$state")
|
||||
if [ -n "$failed" ]; then
|
||||
echo "::error::Cheap check(s) failed on ${PR_SHA}: ${failed}." \
|
||||
"NOT triggering the Buildkite full suite. Push a fix (the 'ready'" \
|
||||
"label re-arms on every push), or re-run the failed check and then" \
|
||||
"re-run this workflow."
|
||||
exit 1
|
||||
fi
|
||||
|
||||
# 'cancelled' counts as pending: wait for a re-run to reach a real verdict
|
||||
# (bounded by MAX_WAIT, then the fail-open below).
|
||||
pending=$(jq '[.[] | select(.status != "completed" or .conclusion == "cancelled")] | length' <<<"$state")
|
||||
missing=$(jq -r --argjson watched "$WATCHED_NAMES" '($watched - map(.name)) | join(", ")' <<<"$state")
|
||||
if [ "$pending" -eq 0 ]; then
|
||||
if [ -z "$missing" ]; then
|
||||
recheck_ready_label
|
||||
echo "All watched cheap checks are green — full suite may proceed."
|
||||
exit 0
|
||||
fi
|
||||
case "$missing" in
|
||||
*pre-commit*)
|
||||
echo "pre-commit run not found for ${PR_SHA} yet; waiting (pre-commit is never path-filtered, so its absence is anomalous)."
|
||||
;;
|
||||
*)
|
||||
if [ "$elapsed" -ge "$GRACE_SECS" ]; then
|
||||
recheck_ready_label
|
||||
echo "::warning::Watched run(s) never appeared for ${PR_SHA}: ${missing} (path-filtered, likely not applicable). Proceeding on the checks that did run."
|
||||
exit 0
|
||||
fi
|
||||
echo "Waiting up to ${GRACE_SECS}s grace for path-filtered run(s) to appear: ${missing}."
|
||||
;;
|
||||
esac
|
||||
fi
|
||||
else
|
||||
api_fails=$(( api_fails + 1 ))
|
||||
echo "::warning::GitHub API error querying workflow runs for ${PR_SHA} (attempt ${api_fails}/3)."
|
||||
if [ "$api_fails" -ge 3 ]; then
|
||||
recheck_ready_label
|
||||
echo "::warning::FAILING OPEN: cannot query GitHub check status — triggering the full suite WITHOUT the cheap-check gate."
|
||||
exit 0
|
||||
fi
|
||||
fi
|
||||
|
||||
if [ "$elapsed" -ge "$MAX_WAIT_SECS" ]; then
|
||||
recheck_ready_label
|
||||
echo "::warning::FAILING OPEN: watched checks still pending after $(( MAX_WAIT_SECS / 60 )) min${missing:+ (never appeared: ${missing})} — triggering the full suite anyway."
|
||||
exit 0
|
||||
fi
|
||||
sleep "$POLL_SECS"
|
||||
done
|
||||
@@ -1,122 +0,0 @@
|
||||
#!/usr/bin/env bash
|
||||
# Self-test for gate_full_suite.sh using a mocked `gh`. No network, runs on
|
||||
# any dev box: bash .github/scripts/test_gate_full_suite.sh
|
||||
set -u
|
||||
here=$(cd "$(dirname "$0")" && pwd)
|
||||
tmp=$(mktemp -d)
|
||||
trap 'rm -rf "$tmp"' EXIT
|
||||
|
||||
# Mock gh. Asserts the exact endpoint (including head_sha) it is called
|
||||
# with — an endpoint typo in the gate script fails the test rather than
|
||||
# silently serving canned data. On the runs endpoint it serves
|
||||
# $MOCK_DIR/response_<call#>.json, sticking on the highest existing file,
|
||||
# and exits 1 if none exist (simulates a GitHub API outage). On the pulls
|
||||
# endpoint it serves $MOCK_DIR/pr.json, defaulting to a 'ready'-labeled PR.
|
||||
cat > "$tmp/gh" <<'EOF'
|
||||
#!/usr/bin/env bash
|
||||
if [ "${1:-}" != "api" ]; then
|
||||
echo "unexpected gh invocation: $*" >> "$MOCK_DIR/endpoint_error"
|
||||
exit 2
|
||||
fi
|
||||
case "${2:-}" in
|
||||
"repos/o/r/actions/runs?head_sha=deadbeef&per_page=100")
|
||||
n=$(( $(cat "$MOCK_DIR/count" 2>/dev/null || echo 0) + 1 ))
|
||||
echo "$n" > "$MOCK_DIR/count"
|
||||
while [ "$n" -gt 0 ]; do
|
||||
if [ -f "$MOCK_DIR/response_$n.json" ]; then
|
||||
cat "$MOCK_DIR/response_$n.json"
|
||||
exit 0
|
||||
fi
|
||||
n=$(( n - 1 ))
|
||||
done
|
||||
echo "api outage" >&2
|
||||
exit 1
|
||||
;;
|
||||
"repos/o/r/pulls/42")
|
||||
if [ -f "$MOCK_DIR/pr.json" ]; then
|
||||
cat "$MOCK_DIR/pr.json"
|
||||
else
|
||||
echo '{"labels": [{"name": "ready"}]}'
|
||||
fi
|
||||
;;
|
||||
*)
|
||||
echo "unexpected gh endpoint: $2" >> "$MOCK_DIR/endpoint_error"
|
||||
exit 2
|
||||
;;
|
||||
esac
|
||||
EOF
|
||||
chmod +x "$tmp/gh"
|
||||
|
||||
PC_OK='{"name": "pre-commit", "id": 1, "status": "completed", "conclusion": "success"}'
|
||||
PC_BAD='{"name": "pre-commit", "id": 1, "status": "completed", "conclusion": "failure"}'
|
||||
PC_PENDING='{"name": "pre-commit", "id": 1, "status": "in_progress", "conclusion": null}'
|
||||
DOCS_OK='{"name": "Deploy Documentation", "id": 2, "status": "completed", "conclusion": "success"}'
|
||||
DOCS_BAD='{"name": "Deploy Documentation", "id": 2, "status": "completed", "conclusion": "failure"}'
|
||||
DOCS_CANCELLED='{"name": "Deploy Documentation", "id": 2, "status": "completed", "conclusion": "cancelled"}'
|
||||
OTHER='{"name": "Trigger Full Suite", "id": 3, "status": "in_progress", "conclusion": null}'
|
||||
NULL_NAME='{"name": null, "id": 4, "status": "completed", "conclusion": "failure"}'
|
||||
PC_OK_RERUN='{"name": "pre-commit", "id": 5, "status": "completed", "conclusion": "success"}'
|
||||
|
||||
fails=0
|
||||
want_log="" # optional: expect() also greps out.log for this regex, then resets
|
||||
pr_json="" # optional: served for the pulls (label re-check) endpoint, then resets
|
||||
raw_body="" # optional: serve responses verbatim instead of wrapping in workflow_runs
|
||||
expect() { # <name> <expected-exit> <response json>...
|
||||
local name=$1 want=$2 dir i=1
|
||||
shift 2
|
||||
dir=$(mktemp -d "$tmp/test_XXXXXX")
|
||||
for body in "$@"; do
|
||||
if [ -n "$raw_body" ]; then
|
||||
printf '%s' "$body" > "$dir/response_$i.json"
|
||||
else
|
||||
printf '{"workflow_runs": [%s]}' "$body" > "$dir/response_$i.json"
|
||||
fi
|
||||
i=$(( i + 1 ))
|
||||
done
|
||||
[ -n "$pr_json" ] && printf '%s' "$pr_json" > "$dir/pr.json"
|
||||
( export PATH="$tmp:$PATH" MOCK_DIR="$dir" PR_SHA=deadbeef PR_NUMBER=42 \
|
||||
GITHUB_REPOSITORY=o/r POLL_SECS=0 GRACE_SECS=1 MAX_WAIT_SECS=3
|
||||
bash "$here/gate_full_suite.sh" > "$dir/out.log" 2>&1 )
|
||||
local rc=$?
|
||||
if [ "$rc" -ne "$want" ]; then
|
||||
echo "FAIL: $name (exit $rc, want $want)"
|
||||
cat "$dir/out.log"
|
||||
fails=1
|
||||
elif [ -f "$dir/endpoint_error" ]; then
|
||||
echo "FAIL: $name (mock gh got an unexpected call)"
|
||||
cat "$dir/endpoint_error"
|
||||
fails=1
|
||||
elif [ -n "$want_log" ] && ! grep -Eq "$want_log" "$dir/out.log"; then
|
||||
echo "FAIL: $name (log does not match: $want_log)"
|
||||
cat "$dir/out.log"
|
||||
fails=1
|
||||
else
|
||||
echo "ok: $name"
|
||||
fi
|
||||
want_log="" pr_json="" raw_body=""
|
||||
}
|
||||
|
||||
expect "both green -> proceed" 0 "$PC_OK, $DOCS_OK, $OTHER, $NULL_NAME"
|
||||
expect "docs build failed -> blocked" 1 "$PC_OK, $DOCS_BAD"
|
||||
expect "pre-commit failed -> blocked" 1 "$PC_BAD"
|
||||
expect "pending then green -> proceed" 0 "$PC_PENDING" "$PC_OK, $DOCS_OK"
|
||||
want_log="never appeared.*Deploy Documentation"
|
||||
expect "docs run absent (path-filtered) -> proceed after grace" 0 "$PC_OK"
|
||||
expect "API outage -> fail open" 0
|
||||
want_log="FAILING OPEN"
|
||||
expect "pending past MAX_WAIT -> fail open" 0 "$PC_PENDING"
|
||||
want_log="FAILING OPEN"
|
||||
expect "unrelated runs only -> no grace, fail open at MAX_WAIT" 0 "$OTHER"
|
||||
expect "cancelled docs then green -> proceed" 0 \
|
||||
"$PC_OK, $DOCS_CANCELLED" "$PC_OK, $DOCS_OK"
|
||||
want_log="FAILING OPEN"
|
||||
expect "cancelled docs forever -> fail open at MAX_WAIT" 0 "$PC_OK, $DOCS_CANCELLED"
|
||||
want_log="FAILING OPEN"
|
||||
expect "pre-commit absent -> no grace, fail open at MAX_WAIT" 0 "$DOCS_OK"
|
||||
expect "duplicate run names -> latest wins" 0 "$PC_BAD, $PC_OK_RERUN, $DOCS_OK"
|
||||
raw_body=1
|
||||
expect "garbage response body -> fail open" 0 "this is not json"
|
||||
pr_json='{"labels": [{"name": "other"}]}'
|
||||
expect "ready label removed mid-gate -> blocked" 1 "$PC_OK, $DOCS_OK"
|
||||
|
||||
exit "$fails"
|
||||
@@ -47,6 +47,3 @@ jobs:
|
||||
- uses: pre-commit/action@v3.0.1
|
||||
with:
|
||||
extra_args: --all-files --hook-stage manual
|
||||
# After pre-commit so a self-test failure cannot mask lint failures.
|
||||
- name: Full-suite gate self-test
|
||||
run: bash .github/scripts/test_gate_full_suite.sh
|
||||
|
||||
@@ -7,7 +7,6 @@ on:
|
||||
permissions:
|
||||
contents: read
|
||||
pull-requests: read
|
||||
actions: read
|
||||
|
||||
concurrency:
|
||||
group: full-suite-${{ github.event.pull_request.number }}
|
||||
@@ -19,8 +18,6 @@ jobs:
|
||||
(github.event.action == 'labeled' && github.event.label.name == 'ready')
|
||||
|| github.event.action == 'synchronize'
|
||||
runs-on: ubuntu-latest
|
||||
# Gate below may wait for cheap checks (up to MAX_WAIT_SECS = 25 min).
|
||||
timeout-minutes: 35
|
||||
steps:
|
||||
- name: Check ready label
|
||||
id: check
|
||||
@@ -52,20 +49,6 @@ jobs:
|
||||
"https://api.buildkite.com/v2/organizations/${{ vars.BUILDKITE_ORG_SLUG }}/pipelines/${{ vars.BUILDKITE_PIPELINE_SLUG }}/builds/${build_num}/cancel"
|
||||
done
|
||||
|
||||
# Checks out the BASE branch (default for pull_request_target), so PR
|
||||
# authors cannot tamper with the gate script.
|
||||
- name: Checkout gate script
|
||||
if: steps.check.outputs.has_ready == 'true'
|
||||
uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
|
||||
|
||||
- name: Wait for pre-commit and docs build
|
||||
if: steps.check.outputs.has_ready == 'true'
|
||||
env:
|
||||
GH_TOKEN: ${{ github.token }}
|
||||
PR_SHA: ${{ github.event.pull_request.head.sha }}
|
||||
PR_NUMBER: ${{ github.event.pull_request.number }}
|
||||
run: bash .github/scripts/gate_full_suite.sh
|
||||
|
||||
- name: Trigger Buildkite Full Suite
|
||||
if: steps.check.outputs.has_ready == 'true'
|
||||
env:
|
||||
|
||||
@@ -0,0 +1,95 @@
|
||||
name: Publish FastVideo Kernel to Hugging Face Hub
|
||||
|
||||
# Version-gated like the PyPI flow (publish-kernel.yml): a push to main only
|
||||
# publishes when the fastvideo-kernel version actually changes. Use
|
||||
# workflow_dispatch as the manual override.
|
||||
on:
|
||||
push:
|
||||
branches:
|
||||
- main
|
||||
paths:
|
||||
- "fastvideo-kernel/pyproject.toml"
|
||||
workflow_dispatch:
|
||||
inputs:
|
||||
branch:
|
||||
description: "Hub branch to upload the build"
|
||||
default: ""
|
||||
required: false
|
||||
|
||||
# A Hub build takes hours; never race two publishes to the same repo.
|
||||
concurrency:
|
||||
group: ${{ github.workflow }}
|
||||
cancel-in-progress: false
|
||||
|
||||
jobs:
|
||||
check-version-change:
|
||||
runs-on: ubuntu-latest
|
||||
outputs:
|
||||
version-changed: ${{ steps.check-version.outputs.changed }}
|
||||
new-version: ${{ steps.check-version.outputs.new-version }}
|
||||
steps:
|
||||
- name: Checkout code
|
||||
uses: actions/checkout@v4
|
||||
with:
|
||||
fetch-depth: 2
|
||||
|
||||
- name: Check if version changed
|
||||
id: check-version
|
||||
run: |
|
||||
cd fastvideo-kernel
|
||||
# Get current commit's version from pyproject.toml
|
||||
# Use ^ to match start of line to avoid matching minimum-version
|
||||
NEW_VERSION=$(grep -oP '^version\s*=\s*"\K[^"]+' pyproject.toml)
|
||||
echo "New version: $NEW_VERSION"
|
||||
|
||||
# Get previous version from git history
|
||||
# Note: git show expects path relative to repo root
|
||||
OLD_VERSION=$(git show HEAD~1:fastvideo-kernel/pyproject.toml | grep -oP '^version\s*=\s*"\K[^"]+' || echo "0.0.0")
|
||||
echo "Old version: $OLD_VERSION"
|
||||
|
||||
if [ "$NEW_VERSION" != "$OLD_VERSION" ]; then
|
||||
echo "Version changed from $OLD_VERSION to $NEW_VERSION"
|
||||
echo "changed=true" >> "$GITHUB_OUTPUT"
|
||||
echo "new-version=$NEW_VERSION" >> "$GITHUB_OUTPUT"
|
||||
else
|
||||
echo "Version did not change"
|
||||
echo "changed=false" >> "$GITHUB_OUTPUT"
|
||||
fi
|
||||
|
||||
build:
|
||||
needs: check-version-change
|
||||
if: ${{ needs.check-version-change.outputs.version-changed == 'true' || github.event_name == 'workflow_dispatch' }}
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
|
||||
- name: Build kernel via HF Jobs and upload to hao-ai-lab/fastvideo-kernel
|
||||
# Pinned to a commit SHA: this third-party action receives HF_TOKEN.
|
||||
uses: huggingface/kernel-builder-job@ee590cee1ff01e922310401f18d5554b8da448ac # main, 2026-06-03
|
||||
with:
|
||||
token: ${{ secrets.HF_TOKEN }}
|
||||
namespace: hao-ai-lab
|
||||
flavor: cpu-xl
|
||||
timeout: "21600"
|
||||
script: |
|
||||
set +x
|
||||
export HF_TOKEN="${{ secrets.HF_TOKEN }}"
|
||||
export GIT_LFS_SKIP_SMUDGE=1
|
||||
|
||||
git clone "${{ github.server_url }}/${{ github.repository }}" FastVideo
|
||||
cd FastVideo
|
||||
git checkout "${{ github.sha }}"
|
||||
git submodule update --init --recursive fastvideo-kernel/include/cutlass fastvideo-kernel/include/tk
|
||||
cd fastvideo-kernel
|
||||
|
||||
UPLOAD_BRANCH="${{ github.event.inputs.branch }}"
|
||||
BRANCH_ARGS=()
|
||||
if [ -n "${UPLOAD_BRANCH}" ]; then
|
||||
BRANCH_ARGS=(--branch "${UPLOAD_BRANCH}")
|
||||
fi
|
||||
|
||||
nix run github:huggingface/kernels#kernel-builder -- build-and-upload \
|
||||
--max-jobs 4 \
|
||||
--cores 8 \
|
||||
--repo-id hao-ai-lab/fastvideo-kernel \
|
||||
"${BRANCH_ARGS[@]}"
|
||||
@@ -99,13 +99,6 @@ status.
|
||||
Full Suite is also path-filtered. It validates broader behavior before Mergify
|
||||
can merge a PR.
|
||||
|
||||
A `ready`-labeled PR does not hit Buildkite immediately:
|
||||
`ci-trigger-full-suite.yml` first runs `.github/scripts/gate_full_suite.sh`,
|
||||
which waits for the cheap Tier-1 checks (pre-commit, docs build) on the PR
|
||||
head. A red cheap check blocks the suite (fail closed; the next push re-arms
|
||||
it), while a GitHub outage or a >25 min wait lets it run anyway (fail open).
|
||||
`/test full` bypasses the gate.
|
||||
|
||||
| Buildkite label | `TEST_TYPE` | Main watched paths |
|
||||
|---|---|---|
|
||||
| SSIM Tests | `ssim` | `fastvideo/**/*.py`, `pyproject.toml`, `docker/Dockerfile` |
|
||||
|
||||
@@ -0,0 +1,37 @@
|
||||
---
|
||||
license: apache-2.0
|
||||
---
|
||||
|
||||
# FastVideo Kernel
|
||||
|
||||
Hub-compatible FastVideo CUDA kernels packaged with Hugging Face `kernel-builder`.
|
||||
|
||||
## Usage
|
||||
|
||||
```python
|
||||
from kernels import get_kernel
|
||||
|
||||
fastvideo_kernel = get_kernel("hao-ai-lab/fastvideo-kernel", version=1)
|
||||
```
|
||||
|
||||
The module exposes:
|
||||
|
||||
- `sta_fwd`
|
||||
- `block_sparse_fwd`
|
||||
- `block_sparse_bwd`
|
||||
- `rms_norm`
|
||||
- `layer_norm`
|
||||
- `int8_quant`
|
||||
- `int8_gemm`
|
||||
|
||||
These APIs are thin wrappers around the native FastVideo kernels. Higher-level
|
||||
fallback logic remains in the regular `fastvideo-kernel` Python package.
|
||||
|
||||
## Requirements
|
||||
|
||||
This first Hub build targets CUDA Hopper (`sm_90a`) because the native attention
|
||||
kernels rely on Hopper/ThunderKittens features.
|
||||
|
||||
## Source
|
||||
|
||||
Source repository: https://github.com/hao-ai-lab/FastVideo/tree/main/fastvideo-kernel
|
||||
@@ -40,6 +40,25 @@ Notes:
|
||||
|
||||
## Installation
|
||||
|
||||
### Hugging Face Kernel Hub
|
||||
|
||||
FastVideo kernels are built and published with the Hugging Face
|
||||
[`kernels`](https://huggingface.co/docs/kernels/index) project. Downstream code
|
||||
loads the precompiled kernels from
|
||||
[`hao-ai-lab/fastvideo-kernel`](https://huggingface.co/hao-ai-lab/fastvideo-kernel)
|
||||
without building locally:
|
||||
|
||||
```python
|
||||
from kernels import get_kernel
|
||||
|
||||
fastvideo_kernel = get_kernel("hao-ai-lab/fastvideo-kernel", version=1)
|
||||
```
|
||||
|
||||
The Hub package exposes the native CUDA entrypoints (`sta_fwd`,
|
||||
`block_sparse_fwd`, `block_sparse_bwd`, `rms_norm`, `layer_norm`, `int8_quant`,
|
||||
and `int8_gemm`). The full FastVideo Python package still provides the
|
||||
high-level fallback logic used during local development.
|
||||
|
||||
### Standard Installation (Local Development)
|
||||
This will automatically detect your GPU architecture. If an NVIDIA Hopper (H100/sm_90a) GPU is detected, ThunderKittens kernels will be enabled. Otherwise, they will be skipped, and the package will use Triton fallbacks at runtime.
|
||||
|
||||
|
||||
@@ -0,0 +1,52 @@
|
||||
[general]
|
||||
name = "fastvideo_kernel"
|
||||
version = 1
|
||||
license = "Apache-2.0"
|
||||
backends = ["cuda"]
|
||||
upstream = "https://github.com/hao-ai-lab/FastVideo.git"
|
||||
source = "https://github.com/hao-ai-lab/FastVideo.git"
|
||||
|
||||
[general.hub]
|
||||
repo-id = "hao-ai-lab/fastvideo-kernel"
|
||||
|
||||
[torch]
|
||||
src = [
|
||||
"torch-ext/torch_binding.cpp",
|
||||
"torch-ext/torch_binding.h",
|
||||
]
|
||||
|
||||
[kernel.fastvideo_kernel]
|
||||
backend = "cuda"
|
||||
depends = ["torch"]
|
||||
include = [
|
||||
"csrc",
|
||||
"csrc/turbodiffusion",
|
||||
"include/cutlass/include",
|
||||
"include/tk/include",
|
||||
"include/tk/prototype",
|
||||
]
|
||||
src = [
|
||||
"csrc/attention/st_attn_h100.cu",
|
||||
"csrc/attention/block_sparse_h100.cu",
|
||||
"csrc/turbodiffusion/gemm/gemm.cu",
|
||||
"csrc/turbodiffusion/norm/rmsnorm.cu",
|
||||
"csrc/turbodiffusion/norm/layernorm.cu",
|
||||
"csrc/turbodiffusion/quant/quant.cu",
|
||||
]
|
||||
# Deliberately Hopper-only (sm_90a): the TK attention kernels here require
|
||||
# Hopper, unlike the multi-arch PyPI wheels (see .github/workflows/publish-kernel.yml).
|
||||
cuda-capabilities = ["9.0a"]
|
||||
cuda-flags = [
|
||||
"-DNDEBUG",
|
||||
"-O3",
|
||||
"-std=c++20",
|
||||
"--use_fast_math",
|
||||
"--expt-extended-lambda",
|
||||
"--expt-relaxed-constexpr",
|
||||
"-Xcompiler=-fno-strict-aliasing",
|
||||
"-Xcompiler=-fPIC",
|
||||
"-DTORCH_COMPILE",
|
||||
"-DKITTENS_HOPPER",
|
||||
"-DFASTVIDEO_KERNEL_DISABLE_PYBIND_REGISTRATION",
|
||||
"-Xptxas=--warn-on-spills",
|
||||
]
|
||||
@@ -71,11 +71,12 @@ void int8_gemm(
|
||||
|
||||
}
|
||||
|
||||
#ifndef FASTVIDEO_KERNEL_DISABLE_PYBIND_REGISTRATION
|
||||
void register_gemm(pybind11::module_ &m) {
|
||||
m.def("gemm_cuda", &int8_gemm);
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -9,10 +9,10 @@
|
||||
#include "norm/layernorm.hpp"
|
||||
|
||||
auto layer_norm(
|
||||
at::Tensor const Input,
|
||||
at::Tensor const Input,
|
||||
float eps,
|
||||
std::optional<at::Tensor const> W,
|
||||
std::optional<at::Tensor const> const B,
|
||||
std::optional<at::Tensor> W,
|
||||
std::optional<at::Tensor> const B,
|
||||
std::optional<at::Tensor> Output
|
||||
) {
|
||||
int64_t const m = Input.size(0);
|
||||
@@ -93,7 +93,8 @@ auto layer_norm(
|
||||
return Output;
|
||||
}
|
||||
|
||||
#ifndef FASTVIDEO_KERNEL_DISABLE_PYBIND_REGISTRATION
|
||||
void register_layer_norm(pybind11::module_ &m) {
|
||||
m.def("layer_norm_cuda", &layer_norm);
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
@@ -75,6 +75,8 @@ auto rms_norm(
|
||||
return Output;
|
||||
}
|
||||
|
||||
#ifndef FASTVIDEO_KERNEL_DISABLE_PYBIND_REGISTRATION
|
||||
void register_rms_norm(pybind11::module_ &m) {
|
||||
m.def("rms_norm_cuda", &rms_norm);
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -70,6 +70,8 @@ auto quant(
|
||||
return std::make_tuple(Output, Output_S);
|
||||
}
|
||||
|
||||
#ifndef FASTVIDEO_KERNEL_DISABLE_PYBIND_REGISTRATION
|
||||
void register_quant(pybind11::module_ &m) {
|
||||
m.def("quant_cuda", &quant);
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,11 @@
|
||||
import torch
|
||||
from kernels import get_kernel
|
||||
|
||||
|
||||
fastvideo_kernel = get_kernel("hao-ai-lab/fastvideo-kernel", version=1)
|
||||
|
||||
x = torch.randn((16, 1024), device="cuda", dtype=torch.float16)
|
||||
weight = torch.ones((1024,), device="cuda", dtype=torch.float16)
|
||||
y = fastvideo_kernel.rms_norm(x, 1e-6, weight)
|
||||
|
||||
print(y.shape)
|
||||
@@ -0,0 +1,18 @@
|
||||
{
|
||||
description = "Flake for FastVideo Kernel Hub builds";
|
||||
|
||||
inputs = {
|
||||
self.submodules = true;
|
||||
kernel-builder.url = "github:huggingface/kernels";
|
||||
};
|
||||
|
||||
outputs =
|
||||
{
|
||||
self,
|
||||
kernel-builder,
|
||||
}:
|
||||
kernel-builder.lib.genKernelFlakeOutputs {
|
||||
inherit self;
|
||||
path = ./.;
|
||||
};
|
||||
}
|
||||
@@ -0,0 +1,122 @@
|
||||
from typing import Optional
|
||||
|
||||
import torch
|
||||
|
||||
from ._ops import ops
|
||||
|
||||
|
||||
def sta_fwd(
|
||||
q: torch.Tensor,
|
||||
k: torch.Tensor,
|
||||
v: torch.Tensor,
|
||||
out: torch.Tensor,
|
||||
kernel_t_size: int,
|
||||
kernel_h_size: int,
|
||||
kernel_w_size: int,
|
||||
text_length: int,
|
||||
process_text: bool,
|
||||
has_text: bool,
|
||||
kernel_aspect_ratio_flag: int,
|
||||
) -> torch.Tensor:
|
||||
return ops.sta_fwd(
|
||||
q,
|
||||
k,
|
||||
v,
|
||||
out,
|
||||
kernel_t_size,
|
||||
kernel_h_size,
|
||||
kernel_w_size,
|
||||
text_length,
|
||||
process_text,
|
||||
has_text,
|
||||
kernel_aspect_ratio_flag,
|
||||
)
|
||||
|
||||
|
||||
def block_sparse_fwd(
|
||||
q: torch.Tensor,
|
||||
k: torch.Tensor,
|
||||
v: torch.Tensor,
|
||||
q2k_block_sparse_index: torch.Tensor,
|
||||
q2k_block_sparse_num: torch.Tensor,
|
||||
kv_block_size: torch.Tensor,
|
||||
) -> tuple[torch.Tensor, torch.Tensor]:
|
||||
out, lse = ops.block_sparse_fwd(
|
||||
q,
|
||||
k,
|
||||
v,
|
||||
q2k_block_sparse_index,
|
||||
q2k_block_sparse_num,
|
||||
kv_block_size,
|
||||
)
|
||||
return out, lse
|
||||
|
||||
|
||||
def block_sparse_bwd(
|
||||
q: torch.Tensor,
|
||||
k: torch.Tensor,
|
||||
v: torch.Tensor,
|
||||
out: torch.Tensor,
|
||||
l_vec: torch.Tensor,
|
||||
out_grad: torch.Tensor,
|
||||
k2q_block_sparse_index: torch.Tensor,
|
||||
k2q_block_sparse_num: torch.Tensor,
|
||||
kv_block_size: torch.Tensor,
|
||||
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
|
||||
q_grad, k_grad, v_grad = ops.block_sparse_bwd(
|
||||
q,
|
||||
k,
|
||||
v,
|
||||
out,
|
||||
l_vec,
|
||||
out_grad,
|
||||
k2q_block_sparse_index,
|
||||
k2q_block_sparse_num,
|
||||
kv_block_size,
|
||||
)
|
||||
return q_grad, k_grad, v_grad
|
||||
|
||||
|
||||
def rms_norm(
|
||||
input: torch.Tensor,
|
||||
eps: float,
|
||||
weight: Optional[torch.Tensor] = None,
|
||||
output: Optional[torch.Tensor] = None,
|
||||
) -> torch.Tensor:
|
||||
return ops.rms_norm_cuda(input, eps, weight, output)
|
||||
|
||||
|
||||
def layer_norm(
|
||||
input: torch.Tensor,
|
||||
eps: float,
|
||||
weight: Optional[torch.Tensor] = None,
|
||||
bias: Optional[torch.Tensor] = None,
|
||||
output: Optional[torch.Tensor] = None,
|
||||
) -> torch.Tensor:
|
||||
return ops.layer_norm_cuda(input, eps, weight, bias, output)
|
||||
|
||||
|
||||
def int8_quant(input: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
|
||||
quantized, scale = ops.quant_cuda(input)
|
||||
return quantized, scale
|
||||
|
||||
|
||||
def int8_gemm(
|
||||
a: torch.Tensor,
|
||||
a_scale: torch.Tensor,
|
||||
b: torch.Tensor,
|
||||
b_scale: torch.Tensor,
|
||||
out: torch.Tensor,
|
||||
) -> torch.Tensor:
|
||||
return ops.gemm_cuda(a, a_scale, b, b_scale, out)
|
||||
|
||||
|
||||
__all__ = [
|
||||
"sta_fwd",
|
||||
"block_sparse_fwd",
|
||||
"block_sparse_bwd",
|
||||
"rms_norm",
|
||||
"layer_norm",
|
||||
"int8_quant",
|
||||
"int8_gemm",
|
||||
]
|
||||
@@ -0,0 +1,99 @@
|
||||
#include <torch/library.h>
|
||||
|
||||
#include "registration.h"
|
||||
#include "torch_binding.h"
|
||||
|
||||
torch::Tensor rms_norm_cuda(
|
||||
torch::Tensor const& input,
|
||||
double eps,
|
||||
std::optional<torch::Tensor> const& weight,
|
||||
std::optional<torch::Tensor> const& output) {
|
||||
std::optional<at::Tensor> mutable_output = output;
|
||||
auto result = rms_norm(input, static_cast<float>(eps), weight, mutable_output);
|
||||
TORCH_CHECK(result.has_value(), "rms_norm_cuda did not produce an output tensor");
|
||||
return result.value();
|
||||
}
|
||||
|
||||
torch::Tensor layer_norm_cuda(
|
||||
torch::Tensor const& input,
|
||||
double eps,
|
||||
std::optional<torch::Tensor> const& weight,
|
||||
std::optional<torch::Tensor> const& bias,
|
||||
std::optional<torch::Tensor> const& output) {
|
||||
auto result = layer_norm(input, static_cast<float>(eps), weight, bias, output);
|
||||
TORCH_CHECK(result.has_value(), "layer_norm_cuda did not produce an output tensor");
|
||||
return result.value();
|
||||
}
|
||||
|
||||
std::tuple<torch::Tensor, torch::Tensor> quant_cuda(torch::Tensor const& input) {
|
||||
std::optional<torch::Tensor> output;
|
||||
std::optional<torch::Tensor> output_scale;
|
||||
auto result = quant(input, output, output_scale);
|
||||
TORCH_CHECK(std::get<0>(result).has_value(), "quant_cuda did not produce an int8 tensor");
|
||||
TORCH_CHECK(std::get<1>(result).has_value(), "quant_cuda did not produce a scale tensor");
|
||||
return std::make_tuple(std::get<0>(result).value(), std::get<1>(result).value());
|
||||
}
|
||||
|
||||
torch::Tensor gemm_cuda(
|
||||
torch::Tensor const& a,
|
||||
torch::Tensor const& a_scale,
|
||||
torch::Tensor const& b,
|
||||
torch::Tensor const& b_scale,
|
||||
torch::Tensor c) {
|
||||
int8_gemm(a, a_scale, b, b_scale, c);
|
||||
return c;
|
||||
}
|
||||
|
||||
TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
|
||||
ops.def(
|
||||
// `out` is written in place and returned, so the return must be declared
|
||||
// as the same alias set — otherwise torch.compile can reorder/eliminate
|
||||
// the mutation and silently produce wrong results.
|
||||
"sta_fwd(Tensor q, Tensor k, Tensor v, Tensor(a!) out, int kernel_t_size, int kernel_h_size, "
|
||||
"int kernel_w_size, int text_length, bool process_text, bool has_text, "
|
||||
"int kernel_aspect_ratio_flag) -> Tensor(a!)");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("sta_fwd", torch::kCUDA, &sta_forward);
|
||||
#endif
|
||||
|
||||
ops.def(
|
||||
"block_sparse_fwd(Tensor q, Tensor k, Tensor v, Tensor q2k_block_sparse_index, "
|
||||
"Tensor q2k_block_sparse_num, Tensor kv_block_size) -> Tensor[]");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("block_sparse_fwd", torch::kCUDA, &block_sparse_attention_forward);
|
||||
#endif
|
||||
|
||||
ops.def(
|
||||
"block_sparse_bwd(Tensor q, Tensor k, Tensor v, Tensor out, Tensor l_vec, Tensor out_grad, "
|
||||
"Tensor k2q_block_sparse_index, Tensor k2q_block_sparse_num, Tensor kv_block_size) -> Tensor[]");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("block_sparse_bwd", torch::kCUDA, &block_sparse_attention_backward);
|
||||
#endif
|
||||
|
||||
// When `output` is provided it is written in place and aliased by the
|
||||
// return value; the schema must say so for torch.compile correctness.
|
||||
ops.def("rms_norm_cuda(Tensor input, float eps, Tensor? weight=None, Tensor(a!)? output=None) -> Tensor(a!)");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("rms_norm_cuda", torch::kCUDA, &rms_norm_cuda);
|
||||
#endif
|
||||
|
||||
ops.def(
|
||||
"layer_norm_cuda(Tensor input, float eps, Tensor? weight=None, Tensor? bias=None, "
|
||||
"Tensor(a!)? output=None) -> Tensor(a!)");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("layer_norm_cuda", torch::kCUDA, &layer_norm_cuda);
|
||||
#endif
|
||||
|
||||
ops.def("quant_cuda(Tensor input) -> (Tensor, Tensor)");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("quant_cuda", torch::kCUDA, &quant_cuda);
|
||||
#endif
|
||||
|
||||
// `c` is the accumulator written in place and returned (same alias set).
|
||||
ops.def("gemm_cuda(Tensor a, Tensor a_scale, Tensor b, Tensor b_scale, Tensor(a!) c) -> Tensor(a!)");
|
||||
#if defined(CUDA_KERNEL)
|
||||
ops.impl("gemm_cuda", torch::kCUDA, &gemm_cuda);
|
||||
#endif
|
||||
}
|
||||
|
||||
REGISTER_EXTENSION(TORCH_EXTENSION_NAME)
|
||||
@@ -0,0 +1,86 @@
|
||||
#pragma once
|
||||
|
||||
#include <optional>
|
||||
#include <tuple>
|
||||
#include <vector>
|
||||
|
||||
#include <torch/torch.h>
|
||||
|
||||
torch::Tensor sta_forward(
|
||||
torch::Tensor q,
|
||||
torch::Tensor k,
|
||||
torch::Tensor v,
|
||||
torch::Tensor o,
|
||||
int kernel_t_size,
|
||||
int kernel_h_size,
|
||||
int kernel_w_size,
|
||||
int text_length,
|
||||
bool process_text,
|
||||
bool has_text,
|
||||
int kernel_aspect_ratio_flag);
|
||||
|
||||
std::vector<torch::Tensor> block_sparse_attention_forward(
|
||||
torch::Tensor q,
|
||||
torch::Tensor k,
|
||||
torch::Tensor v,
|
||||
torch::Tensor q2k_block_sparse_index,
|
||||
torch::Tensor q2k_block_sparse_num,
|
||||
torch::Tensor kv_block_size);
|
||||
|
||||
std::vector<torch::Tensor> block_sparse_attention_backward(
|
||||
torch::Tensor q,
|
||||
torch::Tensor k,
|
||||
torch::Tensor v,
|
||||
torch::Tensor o,
|
||||
torch::Tensor l_vec,
|
||||
torch::Tensor og,
|
||||
torch::Tensor k2q_block_sparse_index,
|
||||
torch::Tensor k2q_block_sparse_num,
|
||||
torch::Tensor kv_block_size);
|
||||
|
||||
std::optional<at::Tensor> rms_norm(
|
||||
at::Tensor const& input,
|
||||
float eps,
|
||||
std::optional<at::Tensor> const& weight,
|
||||
std::optional<at::Tensor>& output);
|
||||
|
||||
std::optional<at::Tensor> layer_norm(
|
||||
at::Tensor const input,
|
||||
float eps,
|
||||
std::optional<at::Tensor> weight,
|
||||
std::optional<at::Tensor> const bias,
|
||||
std::optional<at::Tensor> output);
|
||||
|
||||
std::tuple<std::optional<torch::Tensor>, std::optional<torch::Tensor>> quant(
|
||||
torch::Tensor const& input,
|
||||
std::optional<torch::Tensor>& output,
|
||||
std::optional<torch::Tensor>& output_scale);
|
||||
|
||||
void int8_gemm(
|
||||
at::Tensor const& a,
|
||||
at::Tensor const& a_scale,
|
||||
at::Tensor const& b,
|
||||
at::Tensor const& b_scale,
|
||||
torch::Tensor& c);
|
||||
|
||||
torch::Tensor rms_norm_cuda(
|
||||
torch::Tensor const& input,
|
||||
double eps,
|
||||
std::optional<torch::Tensor> const& weight,
|
||||
std::optional<torch::Tensor> const& output);
|
||||
|
||||
torch::Tensor layer_norm_cuda(
|
||||
torch::Tensor const& input,
|
||||
double eps,
|
||||
std::optional<torch::Tensor> const& weight,
|
||||
std::optional<torch::Tensor> const& bias,
|
||||
std::optional<torch::Tensor> const& output);
|
||||
|
||||
std::tuple<torch::Tensor, torch::Tensor> quant_cuda(torch::Tensor const& input);
|
||||
|
||||
torch::Tensor gemm_cuda(
|
||||
torch::Tensor const& a,
|
||||
torch::Tensor const& a_scale,
|
||||
torch::Tensor const& b,
|
||||
torch::Tensor const& b_scale,
|
||||
torch::Tensor c);
|
||||
@@ -75,25 +75,6 @@ if not is_flash_attn_2_available():
|
||||
|
||||
logger = logging.get_logger(__name__)
|
||||
|
||||
|
||||
def _resolve_torch_dtype(dtype, default: torch.dtype = torch.float32) -> torch.dtype:
|
||||
if isinstance(dtype, torch.dtype):
|
||||
return dtype
|
||||
if isinstance(dtype, str):
|
||||
normalized = dtype.strip().removeprefix("torch.").lower()
|
||||
return {
|
||||
"bfloat16": torch.bfloat16,
|
||||
"bf16": torch.bfloat16,
|
||||
"float16": torch.float16,
|
||||
"fp16": torch.float16,
|
||||
"half": torch.float16,
|
||||
"float32": torch.float32,
|
||||
"fp32": torch.float32,
|
||||
"float": torch.float32,
|
||||
}.get(normalized, default)
|
||||
return default
|
||||
|
||||
|
||||
class Qwen2_5_VLMLP(nn.Module):
|
||||
def __init__(self, config, bias: bool = False):
|
||||
super().__init__()
|
||||
@@ -323,13 +304,10 @@ class Qwen2_5_VisionTransformerPretrainedModel(nn.Module):
|
||||
config_class = Qwen2_5_VLVisionConfig
|
||||
_no_split_modules = ["Qwen2_5_VLVisionBlock"]
|
||||
|
||||
def __init__(self, config, parent_torch_dtype=None) -> None:
|
||||
def __init__(self, config) -> None:
|
||||
super().__init__()
|
||||
|
||||
config_torch_dtype = getattr(config, "torch_dtype", None)
|
||||
self.dtype = _resolve_torch_dtype(
|
||||
config_torch_dtype if config_torch_dtype is not None else parent_torch_dtype
|
||||
)
|
||||
self.dtype = torch.bfloat16 if config.torch_dtype == "bfloat16" else torch.float32
|
||||
|
||||
self.spatial_merge_size = config.spatial_merge_size
|
||||
self.patch_size = config.patch_size
|
||||
@@ -1480,10 +1458,7 @@ class Qwen2_5_VLForConditionalGenerationSimple(nn.Module):
|
||||
super().__init__()
|
||||
config = _flatten_text_config(config)
|
||||
self.config = config
|
||||
self.visual = Qwen2_5_VisionTransformerPretrainedModel(
|
||||
config.vision_config,
|
||||
parent_torch_dtype=getattr(config, "torch_dtype", None),
|
||||
)
|
||||
self.visual = Qwen2_5_VisionTransformerPretrainedModel(config.vision_config)
|
||||
|
||||
self.model = Qwen2_5_VLModel(config)
|
||||
self.vocab_size = config.vocab_size
|
||||
|
||||
@@ -1,58 +0,0 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
"""Guard the Modal FA4 defaults that keep CI lanes on their intended backend.
|
||||
|
||||
Pure text/AST analysis: no fastvideo imports, no torch, no Modal client.
|
||||
"""
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
import ast
|
||||
from pathlib import Path
|
||||
|
||||
REPO_ROOT = Path(__file__).resolve().parents[3]
|
||||
MODAL_ROOT = REPO_ROOT / "fastvideo" / "tests" / "modal"
|
||||
PR_TEST = MODAL_ROOT / "pr_test.py"
|
||||
LAUNCH_L40S_JOB = MODAL_ROOT / "launch_l40s_job.py"
|
||||
SSIM_TEST = MODAL_ROOT / "ssim_test.py"
|
||||
|
||||
|
||||
def _function_strings(path: Path, function_name: str) -> str:
|
||||
source = path.read_text(encoding="utf-8")
|
||||
tree = ast.parse(source)
|
||||
for node in tree.body:
|
||||
if isinstance(node, ast.FunctionDef) and node.name == function_name:
|
||||
return "\n".join(
|
||||
child.value
|
||||
for child in ast.walk(node)
|
||||
if isinstance(child, ast.Constant)
|
||||
and isinstance(child.value, str)
|
||||
)
|
||||
raise AssertionError(f"{function_name} not found in {path}")
|
||||
|
||||
|
||||
def test_generic_l40s_launcher_defaults_fa4_off():
|
||||
source = LAUNCH_L40S_JOB.read_text(encoding="utf-8")
|
||||
assert '"FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "0")' in source
|
||||
|
||||
|
||||
def test_ssim_launcher_keeps_fa4_enabled_by_default():
|
||||
source = SSIM_TEST.read_text(encoding="utf-8")
|
||||
assert '"FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "1")' in source
|
||||
|
||||
|
||||
def test_pr_model_load_and_training_lanes_disable_fa4():
|
||||
lanes = {
|
||||
"run_transformer_tests": "pytest ./fastvideo/tests/transformers -vs",
|
||||
"run_training_tests": "pytest ./fastvideo/tests/training/Vanilla -srP",
|
||||
"run_training_lora_tests": "pytest ./fastvideo/tests/training/lora/test_lora_training.py -srP",
|
||||
"run_training_tests_VSA": "pytest ./fastvideo/tests/training/VSA -srP",
|
||||
"run_distill_dmd_tests": "pytest ./fastvideo/tests/training/distill/test_distill_dmd.py -vs",
|
||||
"run_self_forcing_tests": "pytest ./fastvideo/tests/training/self-forcing/test_self_forcing.py -vs",
|
||||
"run_train_framework_tests": "pytest ./fastvideo/tests/train/models ./fastvideo/tests/train/methods -vs",
|
||||
"seed_grad_norm_references": "pytest ./fastvideo/tests/train/methods -vs -rs",
|
||||
}
|
||||
|
||||
for function_name, pytest_command in lanes.items():
|
||||
function_strings = _function_strings(PR_TEST, function_name)
|
||||
assert "FASTVIDEO_FA4=0" in function_strings
|
||||
assert pytest_command in function_strings
|
||||
@@ -1,96 +0,0 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
|
||||
from types import SimpleNamespace
|
||||
|
||||
import torch
|
||||
|
||||
from fastvideo.models.encoders.qwen2_5_vl_custom import (
|
||||
Qwen2_5_VisionTransformerPretrainedModel,
|
||||
Qwen2_5_VLForConditionalGenerationSimple,
|
||||
)
|
||||
|
||||
|
||||
def _vision_config(torch_dtype=None):
|
||||
return SimpleNamespace(
|
||||
torch_dtype=torch_dtype,
|
||||
spatial_merge_size=2,
|
||||
patch_size=14,
|
||||
fullatt_block_indexes=[],
|
||||
window_size=112,
|
||||
temporal_patch_size=2,
|
||||
in_channels=3,
|
||||
hidden_size=8,
|
||||
num_heads=2,
|
||||
depth=0,
|
||||
_attn_implementation="sdpa",
|
||||
out_hidden_size=8,
|
||||
)
|
||||
|
||||
|
||||
def _full_config(torch_dtype):
|
||||
return SimpleNamespace(
|
||||
vision_config=_vision_config(),
|
||||
torch_dtype=torch_dtype,
|
||||
vocab_size=16,
|
||||
hidden_size=8,
|
||||
num_hidden_layers=0,
|
||||
num_attention_heads=2,
|
||||
pad_token_id=0,
|
||||
_attn_implementation="sdpa",
|
||||
rms_norm_eps=1e-6,
|
||||
max_position_embeddings=32,
|
||||
rope_scaling=None,
|
||||
)
|
||||
|
||||
|
||||
def test_vision_dtype_uses_parent_bfloat16_string_when_vision_dtype_missing():
|
||||
model = Qwen2_5_VisionTransformerPretrainedModel(
|
||||
_vision_config(),
|
||||
parent_torch_dtype="bfloat16",
|
||||
)
|
||||
|
||||
assert model.dtype == torch.bfloat16
|
||||
|
||||
|
||||
def test_vision_dtype_accepts_parent_torch_dtype_object():
|
||||
model = Qwen2_5_VisionTransformerPretrainedModel(
|
||||
_vision_config(),
|
||||
parent_torch_dtype=torch.bfloat16,
|
||||
)
|
||||
|
||||
assert model.dtype == torch.bfloat16
|
||||
|
||||
|
||||
def test_vision_dtype_strips_parent_dtype_string_whitespace():
|
||||
model = Qwen2_5_VisionTransformerPretrainedModel(
|
||||
_vision_config(),
|
||||
parent_torch_dtype=" torch.bfloat16 ",
|
||||
)
|
||||
|
||||
assert model.dtype == torch.bfloat16
|
||||
|
||||
|
||||
def test_vision_dtype_prefers_explicit_vision_dtype_over_parent_dtype():
|
||||
model = Qwen2_5_VisionTransformerPretrainedModel(
|
||||
_vision_config(torch_dtype="float16"),
|
||||
parent_torch_dtype="bfloat16",
|
||||
)
|
||||
|
||||
assert model.dtype == torch.float16
|
||||
|
||||
|
||||
def test_vision_dtype_falls_back_to_float32_for_missing_or_unknown_dtype():
|
||||
missing = Qwen2_5_VisionTransformerPretrainedModel(_vision_config())
|
||||
unknown = Qwen2_5_VisionTransformerPretrainedModel(
|
||||
_vision_config(),
|
||||
parent_torch_dtype="not-a-real-dtype",
|
||||
)
|
||||
|
||||
assert missing.dtype == torch.float32
|
||||
assert unknown.dtype == torch.float32
|
||||
|
||||
|
||||
def test_conditional_generation_passes_parent_dtype_to_visual_tower():
|
||||
model = Qwen2_5_VLForConditionalGenerationSimple(_full_config("torch.bfloat16"))
|
||||
|
||||
assert model.visual.dtype == torch.bfloat16
|
||||
@@ -97,10 +97,9 @@ image = (
|
||||
"TOKENIZERS_PARALLELISM": "false",
|
||||
**({"UV_TORCH_BACKEND": uv_torch_backend_override} if uv_torch_backend_override else {}),
|
||||
"FASTVIDEO_ATTENTION_BACKEND": os.environ.get("FASTVIDEO_ATTENTION_BACKEND", "FLASH_ATTN"),
|
||||
# FA4 is opt-in (FASTVIDEO_FA4). Generic ad hoc jobs should follow the
|
||||
# product default unless a caller opts in through the local env or
|
||||
# --env-vars.
|
||||
"FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "0"),
|
||||
# FA4 is opt-in (FASTVIDEO_FA4); keep CI parity with the seeded
|
||||
# references. Caller override wins.
|
||||
"FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "1"),
|
||||
})
|
||||
)
|
||||
|
||||
|
||||
@@ -73,10 +73,8 @@ ci_env_secret = modal.Secret.from_dict({
|
||||
**({
|
||||
"UV_TORCH_BACKEND": uv_torch_backend_override
|
||||
} if uv_torch_backend_override else {}),
|
||||
# FA4 is opt-in (FASTVIDEO_FA4). Keep the default enabled for
|
||||
# inference/perf parity; model-load and training lanes that do not exercise
|
||||
# FA4 explicitly set FASTVIDEO_FA4=0 in their command strings below.
|
||||
# Caller override wins.
|
||||
# FA4 is opt-in (FASTVIDEO_FA4); CI lanes keep it enabled to match the
|
||||
# SSIM/perf baselines. Caller override wins.
|
||||
"FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "1"),
|
||||
})
|
||||
|
||||
@@ -186,8 +184,7 @@ def run_vae_tests():
|
||||
volumes={"/root/data": model_vol})
|
||||
def run_transformer_tests():
|
||||
run_test(
|
||||
"export HF_HOME='/root/data/.cache' && hf auth login --token $HF_API_KEY && "
|
||||
"FASTVIDEO_FA4=0 pytest ./fastvideo/tests/transformers -vs"
|
||||
"export HF_HOME='/root/data/.cache' && hf auth login --token $HF_API_KEY && pytest ./fastvideo/tests/transformers -vs"
|
||||
)
|
||||
|
||||
|
||||
@@ -200,8 +197,7 @@ def run_transformer_tests():
|
||||
volumes={"/root/data": model_vol})
|
||||
def run_training_tests():
|
||||
run_test(
|
||||
"export HF_HOME='/root/data/.cache' && wandb login $WANDB_API_KEY && "
|
||||
"FASTVIDEO_FA4=0 pytest ./fastvideo/tests/training/Vanilla -srP"
|
||||
"export HF_HOME='/root/data/.cache' && wandb login $WANDB_API_KEY && pytest ./fastvideo/tests/training/Vanilla -srP"
|
||||
)
|
||||
|
||||
|
||||
@@ -214,8 +210,7 @@ def run_training_tests():
|
||||
volumes={"/root/data": model_vol})
|
||||
def run_training_lora_tests():
|
||||
run_test(
|
||||
"export HF_HOME='/root/data/.cache' && wandb login $WANDB_API_KEY && "
|
||||
"FASTVIDEO_FA4=0 pytest ./fastvideo/tests/training/lora/test_lora_training.py -srP"
|
||||
"export HF_HOME='/root/data/.cache' && wandb login $WANDB_API_KEY && pytest ./fastvideo/tests/training/lora/test_lora_training.py -srP"
|
||||
)
|
||||
|
||||
|
||||
@@ -225,7 +220,7 @@ def run_training_lora_tests():
|
||||
secrets=[wandb_secret, ci_env_secret])
|
||||
def run_training_tests_VSA():
|
||||
run_test(
|
||||
"wandb login $WANDB_API_KEY && FASTVIDEO_FA4=0 pytest ./fastvideo/tests/training/VSA -srP"
|
||||
"wandb login $WANDB_API_KEY && pytest ./fastvideo/tests/training/VSA -srP"
|
||||
)
|
||||
|
||||
|
||||
@@ -259,7 +254,7 @@ def run_inference_lora_tests():
|
||||
@app.function(gpu="L40S:2", image=image, timeout=900, secrets=[ci_env_secret])
|
||||
def run_distill_dmd_tests():
|
||||
run_test(
|
||||
"FASTVIDEO_FA4=0 pytest ./fastvideo/tests/training/distill/test_distill_dmd.py -vs")
|
||||
"pytest ./fastvideo/tests/training/distill/test_distill_dmd.py -vs")
|
||||
|
||||
|
||||
@app.function(gpu="L40S:2",
|
||||
@@ -268,8 +263,7 @@ def run_distill_dmd_tests():
|
||||
secrets=[wandb_secret, ci_env_secret])
|
||||
def run_self_forcing_tests():
|
||||
run_test(
|
||||
"wandb login $WANDB_API_KEY && "
|
||||
"FASTVIDEO_FA4=0 pytest ./fastvideo/tests/training/self-forcing/test_self_forcing.py -vs"
|
||||
"wandb login $WANDB_API_KEY && pytest ./fastvideo/tests/training/self-forcing/test_self_forcing.py -vs"
|
||||
)
|
||||
|
||||
|
||||
@@ -329,8 +323,7 @@ def run_dreamverse_app_tests():
|
||||
volumes={"/root/data": model_vol})
|
||||
def run_train_framework_tests():
|
||||
run_test(
|
||||
"export HF_HOME='/root/data/.cache' && hf auth login --token $HF_API_KEY && "
|
||||
"FASTVIDEO_FA4=0 pytest ./fastvideo/tests/train/models ./fastvideo/tests/train/methods -vs"
|
||||
"export HF_HOME='/root/data/.cache' && hf auth login --token $HF_API_KEY && pytest ./fastvideo/tests/train/models ./fastvideo/tests/train/methods -vs"
|
||||
)
|
||||
|
||||
|
||||
@@ -356,8 +349,7 @@ def seed_grad_norm_references():
|
||||
the local command and the ``_DEVICE_MAPPINGS`` table.
|
||||
"""
|
||||
run_test(
|
||||
"export HF_HOME='/root/data/.cache' && hf auth login --token $HF_API_KEY && "
|
||||
"FASTVIDEO_FA4=0 FASTVIDEO_GRADNORM_UPDATE=1 pytest ./fastvideo/tests/train/methods -vs -rs"
|
||||
"export HF_HOME='/root/data/.cache' && hf auth login --token $HF_API_KEY && FASTVIDEO_GRADNORM_UPDATE=1 pytest ./fastvideo/tests/train/methods -vs -rs"
|
||||
)
|
||||
|
||||
|
||||
|
||||
Reference in New Issue
Block a user