[Bench][Linear-Attn] Add Qwen GDN prefill benchmark surface

.github/actions/reclaim-runner-disk/action.yml

@@ -1,24 +1,13 @@
 name: Reclaim runner disk
 description: >
-  Reclaim disk on self-hosted CI runners before heavy steps run. Prunes
-  stale trusted venvs, orphan per-run temp dirs, and aged compile/wheel/pip
-  caches under /home/ci-runner. Fails fast if the volume is still near full
-  after reclaim, so downstream steps report disk pressure clearly instead of
-  a cryptic mkdir error.
+  Trim the persistent bind-mounted /ci-cache (compiled-kernel / triton / pip
+  caches and the tilelang autotuner cache) on self-hosted CI runners. Fails
+  fast if the cache volume is still near full after reclaim, so downstream
+  steps report disk pressure clearly instead of a cryptic mkdir error.
 
 inputs:
-  venv-keep:
-    description: Number of most-recent tileops_ci_venv_* directories to retain.
-    required: false
-    default: "2"
-  temp-orphan-minutes:
-    description: >
-      Age threshold (minutes) for purging orphan gpu-smoke-*/nightly-* temp
-      dirs in RUNNER_TEMP that do not belong to the current run.
-    required: false
-    default: "120"
   cache-age-days:
-    description: Age threshold (days) for trimming compile/wheel/pip cache files.
+    description: Age threshold (days) for trimming cache files.
     required: false
     default: "7"
   min-avail-gib:
@@ -39,12 +28,6 @@ inputs:
       same runner. Cheap disk checks still run on every invocation.
     required: false
     default: "0"
-  current-run-dir:
-    description: >
-      Name of the per-run temp directory that must NOT be pruned. Leave blank
-      if the caller does not create one.
-    required: false
-    default: ""
   cache-dirs:
     description: >
       Newline-separated list of cache directory roots whose files older than
@@ -75,13 +58,6 @@ inputs:
     required: false
     default: |
       /ci-cache/tilelang/autotuner
-  prune-tool-cache:
-    description: >
-      Whether to prune old tileops_ci_venv_* directories under runner.tool_cache.
-      Disable on runners that do not use the per-hash trusted-venv scheme
-      (e.g. nightly's docker-based jobs).
-    required: false
-    default: "true"
   df-paths:
     description: Space-separated paths to pass to `df -h` for before/after logging.
     required: false
@@ -109,20 +85,15 @@ runs:
     - name: Reclaim runner disk
       shell: bash
       env:
-        VENV_KEEP: ${{ inputs.venv-keep }}
-        TEMP_ORPHAN_MINUTES: ${{ inputs.temp-orphan-minutes }}
         CACHE_AGE_DAYS: ${{ inputs.cache-age-days }}
         MIN_AVAIL_GIB: ${{ inputs.min-avail-gib }}
         RECLAIM_BELOW_GIB: ${{ inputs.reclaim-below-gib }}
         CACHE_TRIM_COOLDOWN_MINUTES: ${{ inputs.cache-trim-cooldown-minutes }}
-        CURRENT_RUN_DIR: ${{ inputs.current-run-dir }}
         CACHE_DIRS: ${{ inputs.cache-dirs }}
         ATOMIC_CACHE_DIRS: ${{ inputs.atomic-cache-dirs }}
-        PRUNE_TOOL_CACHE: ${{ inputs.prune-tool-cache }}
         DF_PATHS: ${{ inputs.df-paths }}
         FORCE_RECLAIM: ${{ inputs.force-reclaim }}
         SKIP_ATOMIC_AGE_TRIM: ${{ inputs.skip-atomic-age-trim }}
-        TOOL_CACHE: ${{ runner.tool_cache }}
       run: |
         set -uo pipefail
 
@@ -166,7 +137,7 @@ runs:
 
         # Validate numeric inputs up-front so typos fail with a clear error
         # instead of a cryptic bash arithmetic/find error deeper in the script.
-        for var in VENV_KEEP TEMP_ORPHAN_MINUTES CACHE_AGE_DAYS MIN_AVAIL_GIB RECLAIM_BELOW_GIB CACHE_TRIM_COOLDOWN_MINUTES; do
+        for var in CACHE_AGE_DAYS MIN_AVAIL_GIB RECLAIM_BELOW_GIB CACHE_TRIM_COOLDOWN_MINUTES; do
           val="${!var}"
           if [[ ! "${val}" =~ ^[0-9]+$ ]]; then
             echo "::error::reclaim-runner-disk: input ${var}='${val}' is not a non-negative integer"
@@ -178,7 +149,7 @@ runs:
           read -r -a DF_TARGETS <<< "${DF_PATHS}"
         else
           DF_TARGETS=()
-          for path in "${RUNNER_TEMP:-}" "${TOOL_CACHE:-}" /home/ci-runner; do
+          for path in "${RUNNER_TEMP:-}" /ci-cache; do
             [[ -n "${path}" ]] && DF_TARGETS+=( "${path}" )
           done
         fi
@@ -224,46 +195,7 @@ runs:
           echo "Free space is healthy; skipping reclaim pass"
         fi
 
-        # 1) Prune old trusted venvs: keep N most recent by mtime.
-        if [[ "${SHOULD_RECLAIM}" == "true" && "${PRUNE_TOOL_CACHE}" == "true" && -d "${TOOL_CACHE}" ]]; then
-          DID_RECLAIM="true"
-          mapfile -t venvs < <(find "${TOOL_CACHE}" -maxdepth 1 -type d -name "tileops_ci_venv_*" -printf '%T@ %p\n' 2>/dev/null | sort -rn | cut -d' ' -f2-)
-          if [[ ${#venvs[@]} -gt ${VENV_KEEP} ]]; then
-            for stale in "${venvs[@]:${VENV_KEEP}}"; do
-              echo "Removing stale venv: ${stale}"
-              rm -rf "${stale}" || true
-            done
-          fi
-        fi
-
-        # 2) Purge orphan per-run dirs in RUNNER_TEMP whose newest activity in
-        #    the subtree is older than N minutes, excluding the current run's
-        #    directory if provided. Directory mtime alone is unreliable — it
-        #    only updates when entries are added/removed, so an active
-        #    long-running job whose files are being written (but not the dir
-        #    itself) would be wrongly eligible. Use the newest file mtime
-        #    anywhere in the subtree instead.
-        if [[ "${SHOULD_RECLAIM}" == "true" && -d "${RUNNER_TEMP}" ]]; then
-          DID_RECLAIM="true"
-          now_epoch=$(date +%s)
-          stale_seconds=$(( TEMP_ORPHAN_MINUTES * 60 ))
-          for pattern in 'gpu-smoke-*' 'nightly-*'; do
-            while IFS= read -r candidate; do
-              [[ -n "${candidate}" ]] || continue
-              if [[ -n "${CURRENT_RUN_DIR}" && "$(basename "${candidate}")" == "${CURRENT_RUN_DIR}" ]]; then
-                continue
-              fi
-              newest_mtime=$(find "${candidate}" -printf '%T@\n' 2>/dev/null | awk 'NR == 1 || $1 > max { max = $1 } END { if (NR > 0) print max }')
-              [[ -n "${newest_mtime}" ]] || newest_mtime=$(stat -c %Y "${candidate}" 2>/dev/null || echo 0)
-              if awk "BEGIN { exit !(${now_epoch} - ${newest_mtime} >= ${stale_seconds}) }"; then
-                echo "Removing orphan runtime: ${candidate}"
-                rm -rf "${candidate}" 2>/dev/null || true
-              fi
-            done < <(find "${RUNNER_TEMP}" -mindepth 1 -maxdepth 1 -type d -name "${pattern}" 2>/dev/null)
-          done
-        fi
-
-        # 3) Age-based trim of trusted compile/wheel/pip caches, plus
+        # 1) Age-based trim of the file-granularity cache roots, plus
         #    atomic-granularity trim of atomic cache roots.
         SHOULD_TRIM_CACHE="false"
         if [[ "${SHOULD_RECLAIM}" == "true" && ( -n "${CACHE_DIRS}" || -n "${ATOMIC_CACHE_DIRS}" ) ]]; then

.github/runner/Dockerfile

@@ -150,8 +150,14 @@ ENV AGENT_TOOLSDIRECTORY=/home/ci-runner/runner/_work/_tool \
     PIP_CACHE_DIR=/ci-cache/pip \
     PIP_NO_CACHE_DIR=0
 ARG RUNNER_VERSION=2.334.0
+# The /ci-cache chown below serves only the unmounted image-verification path
+# (`docker run <image> python ...`, see README): ci-runner writes the root-created cache dirs
+# directly. In the mounted CI path the host bind-mount shadows these dirs, so the chown is a
+# no-op and host ownership wins.
 # useradd before WORKDIR: home stays ci-runner-owned, and WORKDIR precedes the relative-path RUN (hadolint).
-RUN useradd -m -s /bin/bash ci-runner
+# Pin UID/GID 1000 so the account deterministically owns the shared /ci-cache bind-mount across
+# hosts (it matches the majority of existing warm-cache entries; host normalization preserves them).
+RUN groupadd -g 1000 ci-runner && useradd -u 1000 -g 1000 -m -s /bin/bash ci-runner
 WORKDIR /home/ci-runner/runner
 RUN mkdir -p /home/ci-runner/runner/_work/_tool \
     && curl -fsSL -o runner.tar.gz \

.github/runner/entrypoint.sh

@@ -1,6 +1,10 @@
 #!/usr/bin/env bash
 set -euo pipefail
 
+# Group-writable (0775 dirs / 0664 files) so compiled kernels in the shared /ci-cache are
+# reusable across runner UIDs that share the ci-runner group.
+umask 002
+
 # Ad-hoc command passthrough: `docker run <image> python ...` / `bash -c ...` runs the
 # given command directly (image verification, smoke tests; see README). With no args the
 # container registers an ephemeral self-hosted runner (below).
@@ -37,4 +41,12 @@ trap cleanup EXIT INT TERM
     --disableupdate
 
 # ── Run one job, then exit ───────────────────────────────────────────────────
-./run.sh
+# Make container stop signals reach the runner so an idle ephemeral runner shuts down promptly
+# instead of waiting out docker stop's grace:
+#   - RUNNER_MANUALLY_TRAP_SIG=1 makes run.sh trap SIGTERM/SIGINT and forward SIGINT to the
+#     listener process group (its default mode does not trap, so the signal is otherwise lost).
+#   - exec makes run.sh PID 1 so docker stop's SIGTERM is delivered to it; a bash wrapper as
+#     PID 1 would ignore an untrapped SIGTERM and stall until SIGKILL.
+# Trade-off: a stop signal cancels an in-progress job rather than draining it.
+export RUNNER_MANUALLY_TRAP_SIG=1
+exec ./run.sh

.github/workflows/gpu-smoke.yml

@@ -404,8 +404,6 @@ jobs:
         with:
           reclaim-below-gib: "10"
           cache-trim-cooldown-minutes: "120"
-          # Run-local venv retired: no tileops_ci_venv_* dirs to prune in the tool cache.
-          prune-tool-cache: "false"
           cache-dirs: |
             /ci-cache/triton
             /ci-cache/pip
@@ -419,15 +417,6 @@ jobs:
           # PR would otherwise reuse.
           skip-atomic-age-trim: "true"
 
-      - name: Use run-local Triton cache
-        run: |
-          set -euo pipefail
-          run_triton_cache="/ci-cache/triton/gpu-smoke-${GITHUB_RUN_ID}-${GITHUB_RUN_ATTEMPT}"
-          mkdir -p "${run_triton_cache}"
-          chmod u+rwx "${run_triton_cache}" || true
-          echo "TRITON_CACHE_DIR=${run_triton_cache}" >> "$GITHUB_ENV"
-          echo "Using TRITON_CACHE_DIR=${run_triton_cache}"
-
       - name: Validate GPU frequency
         run: |
           set -euo pipefail

.github/workflows/nightly.yml

@@ -26,41 +26,15 @@ env:
 
 jobs:
   # =========================================================================
-  # Phase 1 — Warm the kernel compilation cache
-  # =========================================================================
-  cache-warmup:
-    if: ${{ github.repository == 'tile-ai/TileOPs' && (github.event_name == 'schedule' || github.ref == 'refs/heads/main') }}
-    timeout-minutes: 60
-    runs-on: [self-hosted, tile-ops, nightly]
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 0
-
-      - name: Verify nightly runner environment
-        run: scripts/ci/verify_nightly_runner.sh
-        shell: bash
-
-      - name: Install TileOPs (image-baked stack, --no-deps)
-        run: bash scripts/ci/install_tileops.sh
-        shell: bash
-
-      - name: Warm kernel compilation cache
-        continue-on-error: true
-        run: |
-          set -euo pipefail
-          export PYTHONPATH="${GITHUB_WORKSPACE}${PYTHONPATH:+:$PYTHONPATH}"
-          echo "Starting kernel cache warmup..."
-          python3 scripts/warmup_kernel_cache.py --max-workers 64 -n 16
-        shell: bash
-
-  # =========================================================================
-  # Phase 2 — Benchmark (exclusive GPU access for accurate profiling)
+  # Phase 1 — Benchmark (exclusive GPU access for accurate profiling)
+  #
+  # The persistent /ci-cache (TILELANG_CACHE_DIR) carries compiled kernels and
+  # autotuner results across runs; benchmark setup repopulates it on a cold key
+  # (e.g. after a tilelang bump). No separate warmup phase. For a one-off cold
+  # cache, run scripts/warmup_kernel_cache.py manually.
   # =========================================================================
   benchmark:
-    if: ${{ always() && github.repository == 'tile-ai/TileOPs' && (github.event_name == 'schedule' || github.ref == 'refs/heads/main') }}
-    needs: [cache-warmup]
+    if: ${{ github.repository == 'tile-ai/TileOPs' && (github.event_name == 'schedule' || github.ref == 'refs/heads/main') }}
     timeout-minutes: 120
     runs-on: [self-hosted, tile-ops, nightly]
     env:
@@ -135,7 +109,13 @@ jobs:
         shell: bash
 
   # =========================================================================
-  # Phase 3 — Op correctness tests (serial after benchmark for GPU exclusivity)
+  # Phase 2 — Op correctness tests (serial after benchmark for GPU exclusivity)
+  #
+  # Runs the `full` and `nightly` tiers. The `smoke` tier is the PR critical
+  # path and is already validated by every PR and every push-to-main gpu-smoke
+  # (-m "smoke or full"), so re-running it here is pure duplication. `nightly`
+  # is the heavy tier that runs nowhere else; `full` stays as a scheduled
+  # regression / environment-drift net.
   # =========================================================================
   op_test:
     if: ${{ always() && github.repository == 'tile-ai/TileOPs' && (github.event_name == 'schedule' || github.ref == 'refs/heads/main') }}
@@ -170,7 +150,7 @@ jobs:
 
           set -o pipefail
           python3 -m pip install -q -c constraints.txt pytest-timeout
-          python3 -m pytest tests/ -v --tb=short --timeout=900 --timeout-method=thread --junit-xml=test_results.xml | tee tileops_op_test.log
+          python3 -m pytest tests/ -m "full or nightly" -v --tb=short --timeout=900 --timeout-method=thread --junit-xml=test_results.xml | tee tileops_op_test.log
         shell: bash
 
       - name: Download benchmark artifacts
@@ -296,7 +276,7 @@ jobs:
         shell: bash
 
   # =========================================================================
-  # Phase 3b — Publish benchmark data to the nightly-bench orphan branch
+  # Phase 2b — Publish benchmark data to the nightly-bench orphan branch
   # Consumed by tile-ai/TileOPs.github.io to render the docs Benchmarks page.
   # Runs on a github-hosted runner (no GPU needed) and uses the default token
   # to force-push the day's XML snapshot, mirroring the manifest-stats pattern.
@@ -365,7 +345,7 @@ jobs:
           user_email: "41898282+github-actions[bot]@users.noreply.github.com"
 
   # =========================================================================
-  # Phase 4 — Packaging and smoke test
+  # Phase 3 — Packaging and smoke test
   # =========================================================================
   packaging:
     if: ${{ always() && github.repository == 'tile-ai/TileOPs' && (github.event_name == 'schedule' || github.ref == 'refs/heads/main') }}

THIRD_PARTY_NOTICES.md

@@ -0,0 +1,33 @@
+# Third-Party Notices
+
+This repository includes or adapts code from the following MIT-licensed
+projects. The top-level `LICENSE` contains the MIT license text used for this
+repository distribution; the original copyright notices below are retained in
+the relevant source files.
+
+## Qwen FlashQLA
+
+Files under `tileops/kernels/gated_deltanet/gdn_prefill/` implement a TileOps
+version of the CP-split Gated DeltaNet prefill schedule. The schedule-level
+reference for the h-state / corrected-segment-start part of this implementation
+comes from Qwen FlashQLA:
+
+- `__init__.py`
+- `cp_fwd.py`
+- `fused_fwd.py`
+- `prepare_h.py`
+- `tilelang_compat.py`
+- `utils.py`
+
+Original notice:
+
+```text
+Copyright (c) 2026 The Qwen team, Alibaba Group.
+Licensed under the MIT License.
+```
+
+The TileOps versions are not direct wrappers around the FlashQLA kernels. They
+adapt the CP-split scheduling idea into the TileOps operator API, BTHD dispatch,
+TileLang compatibility layer, benchmarking, tests, and local replay/output
+implementation. The utility file keeps the upstream copyright notice present in
+the FlashQLA utility source.