From 5bb2d53ec14f0d802726bc879b994c0286f85415 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Thu, 8 Oct 2026 04:15:24 -0400 Subject: vgsp+guest-diag: kcore-truth pin resolver (pool/pushbuffer/channel ring+userd) + GR ring userd-derived candidates guest_pin_stream.py now resolves translations entirely from guest truth via /proc/kcore + a manual page-table walk (no QMP races): the CE semaphore pool and pushbuffer memory from the UVM rm_mem structures, and each channel's GPFIFO ring + GPPUT from the tracking semaphore's embedding uvm_channel_t (channel_info located by its gpFifoGpuVa value). This removes the watcher pipeline's flakiness on the cudaMalloc stall class (missing pool/pushbuffer pins). gsp_sema_service.c: svc_gr_ring_autopin tries the guest-registered USERD memdesc base as a ring candidate (the fixed FB convention only matches when the guest's VRAM allocator happens to reproduce the old layout), plus an optional (gsp_sema_dbg) one-shot FB scan + raw-word diagnostics when every candidate fails. tools/cuda-ab/p2: self-checking arbitrary-kernel acceptance suite (10 programs: sequential launches, shapes, params, smem/GEMM, streams, malloc churn, printf+error, FMA benchmark, mixed, managed probe) with host oracle expectations and a guest runner. --- tools/cuda-ab/p2/.gitignore | 4 + tools/cuda-ab/p2/README.md | 102 +++++ tools/cuda-ab/p2/build.sh | 39 ++ tools/cuda-ab/p2/expected/SOURCE.txt | 5 + tools/cuda-ab/p2/expected/p2_bench.full.txt | 3 + tools/cuda-ab/p2/expected/p2_bench.txt | 2 + tools/cuda-ab/p2/expected/p2_churn.full.txt | 9 + tools/cuda-ab/p2/expected/p2_churn.txt | 9 + .../cuda-ab/p2/expected/p2_managed_probe.full.txt | 2 + tools/cuda-ab/p2/expected/p2_managed_probe.txt | 2 + tools/cuda-ab/p2/expected/p2_mixed.full.txt | 8 + tools/cuda-ab/p2/expected/p2_mixed.txt | 8 + tools/cuda-ab/p2/expected/p2_params.full.txt | 2 + tools/cuda-ab/p2/expected/p2_params.txt | 2 + tools/cuda-ab/p2/expected/p2_printf_err.full.txt | 10 + tools/cuda-ab/p2/expected/p2_printf_err.txt | 10 + tools/cuda-ab/p2/expected/p2_seq5.full.txt | 7 + tools/cuda-ab/p2/expected/p2_seq5.txt | 7 + tools/cuda-ab/p2/expected/p2_shapes.full.txt | 7 + tools/cuda-ab/p2/expected/p2_shapes.txt | 7 + tools/cuda-ab/p2/expected/p2_smem.full.txt | 3 + tools/cuda-ab/p2/expected/p2_smem.txt | 3 + tools/cuda-ab/p2/expected/p2_streams.full.txt | 3 + tools/cuda-ab/p2/expected/p2_streams.txt | 3 + tools/cuda-ab/p2/oracle.sh | 77 ++++ tools/cuda-ab/p2/run_guest.sh | 161 +++++++ tools/cuda-ab/p2/src/common.h | 66 +++ tools/cuda-ab/p2/src/p2_bench.cu | 95 +++++ tools/cuda-ab/p2/src/p2_churn.cu | 100 +++++ tools/cuda-ab/p2/src/p2_managed_probe.cu | 81 ++++ tools/cuda-ab/p2/src/p2_mixed.cu | 273 ++++++++++++ tools/cuda-ab/p2/src/p2_params.cu | 185 ++++++++ tools/cuda-ab/p2/src/p2_printf_err.cu | 100 +++++ tools/cuda-ab/p2/src/p2_seq5.cu | 127 ++++++ tools/cuda-ab/p2/src/p2_shapes.cu | 111 +++++ tools/cuda-ab/p2/src/p2_smem.cu | 144 +++++++ tools/cuda-ab/p2/src/p2_streams.cu | 109 +++++ tools/guest-diag/guest_pin_stream.py | 474 ++++++++++++++++++++- 38 files changed, 2355 insertions(+), 5 deletions(-) create mode 100644 tools/cuda-ab/p2/.gitignore create mode 100644 tools/cuda-ab/p2/README.md create mode 100755 tools/cuda-ab/p2/build.sh create mode 100644 tools/cuda-ab/p2/expected/SOURCE.txt create mode 100644 tools/cuda-ab/p2/expected/p2_bench.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_bench.txt create mode 100644 tools/cuda-ab/p2/expected/p2_churn.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_churn.txt create mode 100644 tools/cuda-ab/p2/expected/p2_managed_probe.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_managed_probe.txt create mode 100644 tools/cuda-ab/p2/expected/p2_mixed.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_mixed.txt create mode 100644 tools/cuda-ab/p2/expected/p2_params.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_params.txt create mode 100644 tools/cuda-ab/p2/expected/p2_printf_err.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_printf_err.txt create mode 100644 tools/cuda-ab/p2/expected/p2_seq5.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_seq5.txt create mode 100644 tools/cuda-ab/p2/expected/p2_shapes.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_shapes.txt create mode 100644 tools/cuda-ab/p2/expected/p2_smem.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_smem.txt create mode 100644 tools/cuda-ab/p2/expected/p2_streams.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_streams.txt create mode 100755 tools/cuda-ab/p2/oracle.sh create mode 100755 tools/cuda-ab/p2/run_guest.sh create mode 100644 tools/cuda-ab/p2/src/common.h create mode 100644 tools/cuda-ab/p2/src/p2_bench.cu create mode 100644 tools/cuda-ab/p2/src/p2_churn.cu create mode 100644 tools/cuda-ab/p2/src/p2_managed_probe.cu create mode 100644 tools/cuda-ab/p2/src/p2_mixed.cu create mode 100644 tools/cuda-ab/p2/src/p2_params.cu create mode 100644 tools/cuda-ab/p2/src/p2_printf_err.cu create mode 100644 tools/cuda-ab/p2/src/p2_seq5.cu create mode 100644 tools/cuda-ab/p2/src/p2_shapes.cu create mode 100644 tools/cuda-ab/p2/src/p2_smem.cu create mode 100644 tools/cuda-ab/p2/src/p2_streams.cu (limited to 'tools') diff --git a/tools/cuda-ab/p2/.gitignore b/tools/cuda-ab/p2/.gitignore new file mode 100644 index 0000000..57e7c53 --- /dev/null +++ b/tools/cuda-ab/p2/.gitignore @@ -0,0 +1,4 @@ +# Built CUDA binaries (host-built via ./build.sh, pushed to the guest by run_guest.sh) +bin/ +# Guest run outputs (per-run evidence, not references) +guest/ diff --git a/tools/cuda-ab/p2/README.md b/tools/cuda-ab/p2/README.md new file mode 100644 index 0000000..da4e896 --- /dev/null +++ b/tools/cuda-ab/p2/README.md @@ -0,0 +1,102 @@ +# P2 - "arbitrary CUDA kernels" acceptance suite + +A reusable, self-checking acceptance suite for the Open-IOV emulated vGPU. +It extends the single-workload `ab.sh` A/B harness (`victory_mail`, +`vector_add`) to ten independent programs that exercise launch shape, kernel +parameters, shared memory, streams, allocator churn, device printf, error +reporting, a compute-heavy benchmark, and a combined program - each of which +verifies its own results exactly and prints machine-checkable lines. + +``` +p2/ +├── build.sh # nvcc -O2 -arch=sm_75 --cudart=static -> bin/ +├── oracle.sh # host reference run (GPU) or deterministic emit fallback +├── run_guest.sh # push to guest, run with timeout, diff vs expected/, table +├── src/ # common.h + 10 .cu programs +├── bin/ # built binaries (rebuild with ./build.sh) +├── expected/ # .txt stable lines (+ .full.txt, SOURCE.txt) +└── guest/ # per-run guest output (created by run_guest.sh) +``` + +## Programs + +| Program | Covers | Device footprint | +|---|---|---| +| `p2_seq5` | 5 sequential launches in one process on the same buffers, state verified after each | 3 x 4 KB | +| `p2_shapes` | 1x1x1, odd sizes (1013, block 33), non-multiple n (1000/64), 2D grid (40x20), 3D grid (24x8x8) with bounds checks | 8 KB buffer | +| `p2_params` | kernel parameter counts 0..12, types int/float/double/device-pointer | 64 B + 8 B + 4 B | +| `p2_smem` | shared-memory tree reduction (4096 -> 1) and 16x16-tiled 32x32 int GEMM | peak 4 pages | +| `p2_streams` | 3 streams x 2 kernels, per-stream events and result verification | 3 x 4 KB | +| `p2_churn` | 14 cudaMalloc/cudaFree across 6 iterations plus a final fresh alloc | peak 32 KB | +| `p2_printf_err` | device printf (5 deterministic lines) + invalid grid/block launch error reporting + recovery | 4 B | +| `p2_bench` | 8192 threads x 250k correctly-rounded FMAs; fixed verifiable value; TIMING ms/GFLOP/s | 32 KB | +| `p2_mixed` | mail + vadd + int reduction + 2D grid + 2 streams in one process | peak 96 KB | +| `p2_managed_probe` | **expected-failure probe**: cudaMallocManaged outcome recorded, then normal alloc/launch must still work | 4 KB (+4 KB) | + +All arithmetic is exact: integer kernels, fp32 add/mul of small integers +(element comparisons are bit-exact), and correctly-rounded `__fmaf_rn` / +`fmaf`, so `bad=0` is exact and no tolerances are used. + +## Output conventions + +* Compared lines (filtered, sorted): `RESULT`, `CHECK`, `ERR`, `STATUS`, + `P2PRINT`, `PROBE`, `MAIL`. +* Ignored lines: `DEVICE`, `INFO`, `TIMING`, `MEMINFO`. +* Every strict program ends with `STATUS OK` and exits 0; any CUDA failure + prints `ERR cuda=...` + `STATUS FAIL`. + +## Usage + +```bash +./build.sh # build all 10 programs (sm_75, static cudart) + +./oracle.sh # host references -> expected/ + # GPU mode if nvidia-smi works, else + # deterministic --emit-expected fallback + +./run_guest.sh # push + run in the running guest, then + # per-program PASS/FAIL table (exit != 0 on FAIL) +./run_guest.sh p2_seq5 p2_smem # subset +P2_GUEST_TIMEOUT=600 ./run_guest.sh # longer timeout (default 300 s) +``` + +`run_guest.sh` requires the guest to be up (`scripts/vgpu_vm.py exec` must +answer) and `expected/` to exist. It pushes binaries with the same base64 +chunking method as `ab.sh` and never modifies the guest beyond `/tmp/p2_*`. + +## Reference outputs (`expected/`) + +`expected/.txt` holds the stable lines the guest must reproduce; +`expected/.full.txt` keeps the raw oracle output for humans. +`expected/SOURCE.txt` records how the current references were produced. + +* **GPU mode**: `host` runs on bare-metal; full `RUN output` captured. +* **EMIT mode**: when no usable host GPU is available (e.g. the campaign VM + owns the device, `nvidia-smi`/NVML fails), each program is run with + `--emit-expected`, a documented deterministic CPU computation of the + stable lines a correct GPU run must produce. Re-run `./oracle.sh` when a + host GPU is free to replace emit references with captured ones. + +The current `expected/` was generated in **EMIT mode** (see +`expected/SOURCE.txt`) because the host GPU was claimed by the running +guest at build time. + +## Bounce-window budget + +The emulated GR path relocates guest data pointers into bounce windows +(100 pages total, 112-page win0 shadow, 32 windows, ~33 pages per pointer +observed working). This suite stays well inside those caps: largest single +pointer 32 KB (8 pages), largest simultaneous footprint 96 KB (24 pages, +`p2_mixed` vadd phase), total live allocations per program <= 4. No managed +memory is used except in `p2_managed_probe`, which is informational. + +## Special cases + +* `p2_managed_probe` is the one deliberate expected-failure program: UVM + managed memory is a known unsupported path. `run_guest.sh` marks it + `INFO` and does not diff its content, but requires it to finish with + `STATUS OK` (clean managed-alloc failure, normal alloc still working). A + hang is caught by the timeout and fails the suite. +* `p2_printf_err` can legitimately be the most fragile test if the emulated + path does not support device printf buffering; a timeout/`FAIL` there is a + finding, not a harness bug. diff --git a/tools/cuda-ab/p2/build.sh b/tools/cuda-ab/p2/build.sh new file mode 100755 index 0000000..867182c --- /dev/null +++ b/tools/cuda-ab/p2/build.sh @@ -0,0 +1,39 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# +# p2/build.sh - build the P2 "arbitrary CUDA kernels" acceptance programs. +# +# Target is sm_75 (Turing TU104) to match the project's GPU. The CUDA +# runtime is linked statically (--cudart=static) so the guest needs only the +# driver's libcuda.so, exactly like ab.sh/victory_mail. +# +# Usage: ./build.sh [program ...] (default: all programs) +set -e + +HERE="$(cd "$(dirname "$0")" && pwd)" +NVCC="${NVCC:-nvcc}" +FLAGS="${P2_NVCC_FLAGS:--O2 -arch=sm_75 --cudart=static}" + +command -v "$NVCC" >/dev/null || { echo "nvcc not found"; exit 1; } + +mkdir -p "$HERE/bin" + +build_one() { + local src="$HERE/src/$1.cu" + local bin="$HERE/bin/$1" + + [ -f "$src" ] || { echo "[build] missing source: $src" >&2; return 1; } + echo "[build] $1" + # shellcheck disable=SC2086 + "$NVCC" $FLAGS -I"$HERE/src" -o "$bin" "$src" +} + +if [ "$#" -gt 0 ]; then + for p in "$@"; do build_one "$p"; done +else + for src in "$HERE"/src/p2_*.cu; do + build_one "$(basename "$src" .cu)" + done +fi + +echo "[build] OK: $(ls "$HERE"/bin/p2_* 2>/dev/null | wc -l) program(s) in $HERE/bin" diff --git a/tools/cuda-ab/p2/expected/SOURCE.txt b/tools/cuda-ab/p2/expected/SOURCE.txt new file mode 100644 index 0000000..a138e80 --- /dev/null +++ b/tools/cuda-ab/p2/expected/SOURCE.txt @@ -0,0 +1,5 @@ +mode=emit +date=2026-10-08T07:01:11Z +host=majik-envy +deterministic CPU computation via --emit-expected +(no usable host GPU; rerun ./oracle.sh when one is available) diff --git a/tools/cuda-ab/p2/expected/p2_bench.full.txt b/tools/cuda-ab/p2/expected/p2_bench.full.txt new file mode 100644 index 0000000..47a9851 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_bench.full.txt @@ -0,0 +1,3 @@ +RESULT bench threads=8192 iters=250000 bad=0 value=0.00999368541 +TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU) +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_bench.txt b/tools/cuda-ab/p2/expected/p2_bench.txt new file mode 100644 index 0000000..525eb2a --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_bench.txt @@ -0,0 +1,2 @@ +RESULT bench threads=8192 iters=250000 bad=0 value=0.00999368541 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_churn.full.txt b/tools/cuda-ab/p2/expected/p2_churn.full.txt new file mode 100644 index 0000000..a9bd20d --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_churn.full.txt @@ -0,0 +1,9 @@ +RESULT churn iter=0 nbuf=1 n=2048 bad=0 +RESULT churn iter=1 nbuf=2 n=2048 bad=0 +RESULT churn iter=2 nbuf=3 n=2048 bad=0 +RESULT churn iter=3 nbuf=4 n=2048 bad=0 +RESULT churn iter=4 nbuf=1 n=2048 bad=0 +RESULT churn iter=5 nbuf=2 n=2048 bad=0 +RESULT churn final_buf=1 n=2048 bad=0 +RESULT churn iters=6 allocs=14 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_churn.txt b/tools/cuda-ab/p2/expected/p2_churn.txt new file mode 100644 index 0000000..a9bd20d --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_churn.txt @@ -0,0 +1,9 @@ +RESULT churn iter=0 nbuf=1 n=2048 bad=0 +RESULT churn iter=1 nbuf=2 n=2048 bad=0 +RESULT churn iter=2 nbuf=3 n=2048 bad=0 +RESULT churn iter=3 nbuf=4 n=2048 bad=0 +RESULT churn iter=4 nbuf=1 n=2048 bad=0 +RESULT churn iter=5 nbuf=2 n=2048 bad=0 +RESULT churn final_buf=1 n=2048 bad=0 +RESULT churn iters=6 allocs=14 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_managed_probe.full.txt b/tools/cuda-ab/p2/expected/p2_managed_probe.full.txt new file mode 100644 index 0000000..b5d6afc --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_managed_probe.full.txt @@ -0,0 +1,2 @@ +PROBE managed state=no-oracle-gpu note=emit-mode +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_managed_probe.txt b/tools/cuda-ab/p2/expected/p2_managed_probe.txt new file mode 100644 index 0000000..b5d6afc --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_managed_probe.txt @@ -0,0 +1,2 @@ +PROBE managed state=no-oracle-gpu note=emit-mode +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_mixed.full.txt b/tools/cuda-ab/p2/expected/p2_mixed.full.txt new file mode 100644 index 0000000..483d2db --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_mixed.full.txt @@ -0,0 +1,8 @@ +MAIL the little gpu man says p2 mixed +RESULT mixed mail bad=0 +RESULT mixed vadd n=8192 bad=0 +RESULT mixed red n=4096 sum=8386560 bad=0 +RESULT mixed grid2d nx=40 ny=20 bad=0 +RESULT mixed stream nstreams=2 n=1024 bad=0 +RESULT mixed total bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_mixed.txt b/tools/cuda-ab/p2/expected/p2_mixed.txt new file mode 100644 index 0000000..483d2db --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_mixed.txt @@ -0,0 +1,8 @@ +MAIL the little gpu man says p2 mixed +RESULT mixed mail bad=0 +RESULT mixed vadd n=8192 bad=0 +RESULT mixed red n=4096 sum=8386560 bad=0 +RESULT mixed grid2d nx=40 ny=20 bad=0 +RESULT mixed stream nstreams=2 n=1024 bad=0 +RESULT mixed total bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_params.full.txt b/tools/cuda-ab/p2/expected/p2_params.full.txt new file mode 100644 index 0000000..7833c07 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_params.full.txt @@ -0,0 +1,2 @@ +RESULT params kernels=13 nparam_max=12 bad=0 checksum=395 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_params.txt b/tools/cuda-ab/p2/expected/p2_params.txt new file mode 100644 index 0000000..7833c07 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_params.txt @@ -0,0 +1,2 @@ +RESULT params kernels=13 nparam_max=12 bad=0 checksum=395 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.full.txt b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt new file mode 100644 index 0000000..e2bf8a4 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt @@ -0,0 +1,10 @@ +P2PRINT hello i=0 +P2PRINT hello i=1 +P2PRINT hello i=2 +P2PRINT hello i=3 +P2PRINT second c=Z v=7 +ERR launch grid0 rc=9 name=invalid configuration argument +ERR launch block0 rejected=1 +ERR cleared rc=0 +RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.txt b/tools/cuda-ab/p2/expected/p2_printf_err.txt new file mode 100644 index 0000000..e2bf8a4 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_printf_err.txt @@ -0,0 +1,10 @@ +P2PRINT hello i=0 +P2PRINT hello i=1 +P2PRINT hello i=2 +P2PRINT hello i=3 +P2PRINT second c=Z v=7 +ERR launch grid0 rc=9 name=invalid configuration argument +ERR launch block0 rejected=1 +ERR cleared rc=0 +RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_seq5.full.txt b/tools/cuda-ab/p2/expected/p2_seq5.full.txt new file mode 100644 index 0000000..0823ab1 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_seq5.full.txt @@ -0,0 +1,7 @@ +RESULT seq step=1 n=1024 bad=0 +RESULT seq step=2 n=1024 bad=0 +RESULT seq step=3 n=1024 bad=0 +RESULT seq step=4 n=1024 bad=0 +RESULT seq step=5 n=1024 bad=0 +RESULT seq launches=5 n=1024 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_seq5.txt b/tools/cuda-ab/p2/expected/p2_seq5.txt new file mode 100644 index 0000000..0823ab1 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_seq5.txt @@ -0,0 +1,7 @@ +RESULT seq step=1 n=1024 bad=0 +RESULT seq step=2 n=1024 bad=0 +RESULT seq step=3 n=1024 bad=0 +RESULT seq step=4 n=1024 bad=0 +RESULT seq step=5 n=1024 bad=0 +RESULT seq launches=5 n=1024 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_shapes.full.txt b/tools/cuda-ab/p2/expected/p2_shapes.full.txt new file mode 100644 index 0000000..36e4e3d --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_shapes.full.txt @@ -0,0 +1,7 @@ +RESULT shapes case=1 n=1 bad=0 +RESULT shapes case=2 n=1013 bad=0 +RESULT shapes case=3 n=1000 bad=0 +RESULT shapes case=4 n=800 bad=0 +RESULT shapes case=5 n=1536 bad=0 +RESULT shapes cases=5 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_shapes.txt b/tools/cuda-ab/p2/expected/p2_shapes.txt new file mode 100644 index 0000000..36e4e3d --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_shapes.txt @@ -0,0 +1,7 @@ +RESULT shapes case=1 n=1 bad=0 +RESULT shapes case=2 n=1013 bad=0 +RESULT shapes case=3 n=1000 bad=0 +RESULT shapes case=4 n=800 bad=0 +RESULT shapes case=5 n=1536 bad=0 +RESULT shapes cases=5 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_smem.full.txt b/tools/cuda-ab/p2/expected/p2_smem.full.txt new file mode 100644 index 0000000..48c8f60 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_smem.full.txt @@ -0,0 +1,3 @@ +RESULT smem reduce n=4096 sum=8386560 bad=0 +RESULT smem gemm m=32 n=32 k=32 checksum=65410 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_smem.txt b/tools/cuda-ab/p2/expected/p2_smem.txt new file mode 100644 index 0000000..48c8f60 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_smem.txt @@ -0,0 +1,3 @@ +RESULT smem reduce n=4096 sum=8386560 bad=0 +RESULT smem gemm m=32 n=32 k=32 checksum=65410 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_streams.full.txt b/tools/cuda-ab/p2/expected/p2_streams.full.txt new file mode 100644 index 0000000..6c60041 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_streams.full.txt @@ -0,0 +1,3 @@ +RESULT streams nstreams=3 launches=6 n=1024 bad=0 +RESULT streams events=3 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_streams.txt b/tools/cuda-ab/p2/expected/p2_streams.txt new file mode 100644 index 0000000..6c60041 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_streams.txt @@ -0,0 +1,3 @@ +RESULT streams nstreams=3 launches=6 n=1024 bad=0 +RESULT streams events=3 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/oracle.sh b/tools/cuda-ab/p2/oracle.sh new file mode 100755 index 0000000..80cff94 --- /dev/null +++ b/tools/cuda-ab/p2/oracle.sh @@ -0,0 +1,77 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# +# p2/oracle.sh - build the P2 programs and record host-side reference output +# into p2/expected/.txt. +# +# Two modes: +# GPU mode - nvidia-smi works: each program is run bare-metal; the stable +# lines (RESULT/CHECK/ERR/STATUS/P2PRINT/PROBE/MAIL) are +# filtered into expected/.txt and the full output is +# kept as expected/.full.txt. +# EMIT mode - no usable host GPU (e.g. the device is owned by a running +# guest): each program is run with --emit-expected, a +# deterministic CPU-only computation of the stable lines a +# correct GPU run must produce. This fallback is what +# distinguishes "oracle not available" from "oracle passed". +# +# The generation mode is recorded in expected/SOURCE.txt. +# +# Usage: ./oracle.sh +set -e + +HERE="$(cd "$(dirname "$0")" && pwd)" +EXP="$HERE/expected" +FILTER='^(RESULT|CHECK|ERR|STATUS|P2PRINT|PROBE|MAIL)' + +"$HERE/build.sh" +mkdir -p "$EXP" + +mode="emit" +if nvidia-smi -L >/dev/null 2>&1; then + mode="gpu" +fi + +FAIL=0 +for bin in "$HERE"/bin/p2_*; do + name="$(basename "$bin")" + out="$EXP/$name.full.txt" + + if [ "$mode" = "gpu" ]; then + echo "[oracle] $name (bare-metal GPU)" + if ! timeout 300 "$bin" > "$out" 2>&1; then + echo "[oracle] FAIL: $name exited non-zero" >&2 + FAIL=1 + continue + fi + if ! grep -q '^STATUS OK$' "$out"; then + echo "[oracle] FAIL: $name did not print STATUS OK" >&2 + FAIL=1 + continue + fi + grep -E "$FILTER" "$out" > "$EXP/$name.txt" + else + echo "[oracle] $name (no host GPU: deterministic emit)" + "$bin" --emit-expected > "$out" 2>&1 + grep -E "$FILTER" "$out" > "$EXP/$name.txt" + fi +done + +if [ "$FAIL" -ne 0 ]; then + echo "[oracle] FAILED - expected/ is incomplete" >&2 + exit 1 +fi + +{ + echo "mode=$mode" + echo "date=$(date -u +%Y-%m-%dT%H:%M:%SZ)" + echo "host=$(hostname)" + if [ "$mode" = "gpu" ]; then + nvidia-smi --query-gpu=name,driver_version --format=csv,noheader | head -1 + else + echo "deterministic CPU computation via --emit-expected" + echo "(no usable host GPU; rerun ./oracle.sh when one is available)" + fi +} > "$EXP/SOURCE.txt" + +echo "[oracle] done ($mode) -> $EXP" diff --git a/tools/cuda-ab/p2/run_guest.sh b/tools/cuda-ab/p2/run_guest.sh new file mode 100755 index 0000000..0290df1 --- /dev/null +++ b/tools/cuda-ab/p2/run_guest.sh @@ -0,0 +1,161 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# +# p2/run_guest.sh - push every P2 program into the running guest, execute it +# with a timeout, and compare its stable output lines against the recorded +# host-oracle reference in p2/expected/.txt. +# +# Prerequisites: +# * p2/expected/.txt exists (run p2/oracle.sh first) +# * the guest is running and reachable through scripts/vgpu_vm.py +# +# This talks to the live VM; run it only when no campaign cycle is using it. +# +# Usage: +# ./run_guest.sh # all programs +# ./run_guest.sh p2_seq5 ... # selected programs +# P2_GUEST_TIMEOUT=600 ./run_guest.sh +# +# Exit status: 0 only if every strict program PASSes (the managed-memory +# probe is informational: it must complete with STATUS OK, its content is +# recorded but not diffed). +set -u + +HERE="$(cd "$(dirname "$0")" && pwd)" +REPO="$(cd "$HERE/../.." && pwd)" +VMEXEC="${P2_VMEXEC:-$REPO/scripts/vgpu_vm.py}" +FILTER='^(RESULT|CHECK|ERR|STATUS|P2PRINT|PROBE|MAIL)' +TIMEOUT="${P2_GUEST_TIMEOUT:-300}" +INFO_PROGS=" p2_managed_probe " +GUEST_OUT="$HERE/guest" + +mkdir -p "$GUEST_OUT" + +PROGS=() +if [ "$#" -gt 0 ]; then + PROGS=("$@") +else + for b in "$HERE"/bin/p2_*; do + PROGS+=("$(basename "$b")") + done +fi + +if [ ! -x "$VMEXEC" ] && [ ! -f "$VMEXEC" ]; then + echo "ERROR: VM executor not found: $VMEXEC" >&2 + exit 2 +fi + +# Build any missing binaries first (bin/ may be absent in a fresh checkout). +need_build=0 +for name in "${PROGS[@]}"; do + [ -x "$HERE/bin/$name" ] || need_build=1 +done +if [ "$need_build" -eq 1 ]; then + "$HERE/build.sh" || exit 2 +fi + +for name in "${PROGS[@]}"; do + if [ ! -x "$HERE/bin/$name" ]; then + echo "ERROR: missing binary bin/$name - run p2/build.sh" >&2 + exit 2 + fi + if [ ! -f "$HERE/expected/$name.txt" ]; then + echo "ERROR: missing expected/$name.txt - run p2/oracle.sh" >&2 + exit 2 + fi +done + +ping="$("$VMEXEC" exec "echo VMPING" 2>&1 || true)" +if [[ "$ping" != *VMPING* ]]; then + echo "ERROR: guest not reachable via $VMEXEC (is the VM running?)" >&2 + exit 2 +fi + +push_bin() { + local name="$1" b64 total off size=70000 chunk out + + b64="$(base64 -w0 "$HERE/bin/$name")" + total=${#b64} + off=0 + "$VMEXEC" exec "rm -f /tmp/$name.b64" >/dev/null 2>&1 + while [ "$off" -lt "$total" ]; do + chunk="${b64:$off:$size}" + "$VMEXEC" exec "echo -n '$chunk' >> /tmp/$name.b64" >/dev/null 2>&1 + off=$((off + size)) + done + out="$("$VMEXEC" exec "base64 -d /tmp/$name.b64 > /tmp/$name && chmod +x /tmp/$name && rm -f /tmp/$name.b64 && echo PUSH_OK" 2>&1)" + [[ "$out" == *PUSH_OK* ]] +} + +declare -a R_NAME R_STATUS R_NOTE + +echo "== P2 guest acceptance: $(date -u +%Y-%m-%dT%H:%M:%SZ) ==" +echo "== timeout=${TIMEOUT}s, expected dir: $HERE/expected ==" + +for name in "${PROGS[@]}"; do + gs="$GUEST_OUT/$name.txt" + fail=0 + note="" + + if ! push_bin "$name"; then + R_NAME+=("$name"); R_STATUS+=("FAIL"); R_NOTE+=("push failed") + continue + fi + + "$VMEXEC" exec --timeout "$TIMEOUT" "/tmp/$name" > "$gs" 2>&1 + rc=$? + + if [[ " $INFO_PROGS " == *" $name "* ]]; then + if [ "$rc" -eq 0 ] && grep -q '^STATUS OK$' "$gs" && grep -q '^PROBE ' "$gs"; then + note="$(grep -m1 '^PROBE ' "$gs" | cut -c1-70)" + R_NAME+=("$name"); R_STATUS+=("INFO"); R_NOTE+=("$note") + else + R_NAME+=("$name"); R_STATUS+=("FAIL") + R_NOTE+=("expected-failure probe did not complete cleanly (rc=$rc)") + fi + continue + fi + + if [ "$rc" -ne 0 ]; then + fail=1 + note="guest exit=$rc" + fi + if ! grep -q '^STATUS OK$' "$gs"; then + fail=1 + note="${note:+$note; }no STATUS OK" + fi + + grep -E "$FILTER" "$gs" | LC_ALL=C sort > "$gs.filtered" + grep -E "$FILTER" "$HERE/expected/$name.txt" | LC_ALL=C sort > "$gs.expected" + + if ! cmp -s "$gs.filtered" "$gs.expected"; then + fail=1 + note="${note:+$note; }stable-line mismatch" + if [ ! -s "$gs.filtered" ]; then + note="${note:+$note; }no stable lines (timeout/hang?)" + fi + fi + + if [ "$fail" -eq 0 ]; then + R_NAME+=("$name"); R_STATUS+=("PASS"); R_NOTE+=("") + else + R_NAME+=("$name"); R_STATUS+=("FAIL"); R_NOTE+=("$note") + fi +done + +echo +printf '%-22s %-6s %s\n' "PROGRAM" "RESULT" "NOTES" +printf '%-22s %-6s %s\n' "----------------------" "------" "-----" +any_fail=0 +for i in "${!R_NAME[@]}"; do + printf '%-22s %-6s %s\n' "${R_NAME[$i]}" "${R_STATUS[$i]}" "${R_NOTE[$i]}" + [ "${R_STATUS[$i]}" = "FAIL" ] && any_fail=1 +done +echo + +if [ "$any_fail" -ne 0 ]; then + echo "P2: FAIL (guest output in $GUEST_OUT, diffs: *.filtered vs *.expected)" + exit 1 +fi +echo "P2: PASS ($((${#R_NAME[@]})) programs; guest output in $GUEST_OUT)" +exit 0 diff --git a/tools/cuda-ab/p2/src/common.h b/tools/cuda-ab/p2/src/common.h new file mode 100644 index 0000000..e1e3137 --- /dev/null +++ b/tools/cuda-ab/p2/src/common.h @@ -0,0 +1,66 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * common.h - shared conventions for the P2 "arbitrary CUDA kernels" suite. + * + * Line conventions (p2/run_guest.sh filters on these): + * + * compared (stable) RESULT / CHECK / ERR / STATUS / P2PRINT / PROBE / MAIL + * informational DEVICE / INFO / TIMING / MEMINFO (ignored by compare) + * + * Every program accepts --emit-expected: it prints the deterministic + * machine-checkable lines a correct GPU run must produce, computed on the + * CPU only, without touching CUDA. p2/oracle.sh uses this fallback when the + * host has no usable GPU (e.g. the device is owned by a running guest), so + * the expected/ references are a documented deterministic computation + * rather than a captured bare-metal run. + * + * All verification in this suite is exact (integers, or bit-identical + * float/integer arithmetic and correctly-rounded fmaf), so `bad=0` is + * meaningful and no tolerances are used. + */ +#ifndef P2_COMMON_H +#define P2_COMMON_H + +#include +#include +#include +#include +#include + +#define P2_CHECK(call) \ + do { \ + cudaError_t e_ = (call); \ + if (e_ != cudaSuccess) { \ + printf("ERR cuda=%s code=%d line=%d\n", \ + cudaGetErrorString(e_), (int)e_, __LINE__); \ + printf("STATUS FAIL\n"); \ + fflush(stdout); \ + return 1; \ + } \ + } while (0) + +static inline int p2_emit_mode(int argc, char **argv) +{ + return (argc > 1 && strcmp(argv[1], "--emit-expected") == 0) ? 1 : 0; +} + +static inline void p2_device_info(void) +{ + cudaDeviceProp prop; + + if (cudaGetDeviceProperties(&prop, 0) == cudaSuccess) + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem); + else + printf("DEVICE unavailable\n"); +} + +static inline int p2_status(int bad) +{ + printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL"); + fflush(stdout); + return bad == 0 ? 0 : 1; +} + +#endif /* P2_COMMON_H */ diff --git a/tools/cuda-ab/p2/src/p2_bench.cu b/tools/cuda-ab/p2/src/p2_bench.cu new file mode 100644 index 0000000..2d1db09 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_bench.cu @@ -0,0 +1,95 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_bench - compute-heavy FMA benchmark with a fixed, verifiable result. + * + * 8192 threads each run a 250,000-iteration serial chain of correctly + * rounded FMAs (__fmaf_rn on device, fmaf on host), producing 2.05e9 FMAs. + * Both implementations are IEEE-correctly-rounded, so the per-thread final + * value must match the host exactly (bad=0); the same value is printed for + * the A/B comparison. A TIMING line (ignored by the harness comparison) + * reports wall-clock ms and GFLOP/s for the performance demonstration. + * Device memory: 32 KB output (8 pages). + */ +#include "common.h" +#include + +#define BENCH_THREADS 8192 +#define BENCH_ITERS 250000 +#define BENCH_A 0.9999f +#define BENCH_B 1.0e-6f + +__global__ void fma_chain(float *out, int iters, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + float x = 0.0f; + + if (i >= n) + return; + for (int t = 0; t < iters; t++) + x = __fmaf_rn(x, BENCH_A, BENCH_B); + out[i] = x; +} + +static float host_chain(int iters) +{ + float x = 0.0f; + + for (int t = 0; t < iters; t++) + x = fmaf(x, BENCH_A, BENCH_B); + return x; +} + +int main(int argc, char **argv) +{ + float ref = host_chain(BENCH_ITERS); + + if (p2_emit_mode(argc, argv)) { + printf("RESULT bench threads=%d iters=%d bad=0 value=%.9g\n", + BENCH_THREADS, BENCH_ITERS, (double)ref); + printf("TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU)\n"); + return p2_status(0); + } + + p2_device_info(); + + float *d_out = NULL, *got = (float *)malloc(sizeof(float) * BENCH_THREADS); + + if (!got) + return 1; + P2_CHECK(cudaMalloc(&d_out, sizeof(float) * BENCH_THREADS)); + + cudaEvent_t t0, t1; + float ms = 0.0f; + + P2_CHECK(cudaEventCreate(&t0)); + P2_CHECK(cudaEventCreate(&t1)); + P2_CHECK(cudaEventRecord(t0)); + fma_chain<<>>(d_out, BENCH_ITERS, + BENCH_THREADS); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaEventRecord(t1)); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaEventElapsedTime(&ms, t0, t1)); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(float) * BENCH_THREADS, + cudaMemcpyDeviceToHost)); + + int bad = 0; + + for (int i = 0; i < BENCH_THREADS; i++) + if (got[i] != ref) + bad++; + printf("RESULT bench threads=%d iters=%d bad=%d value=%.9g\n", + BENCH_THREADS, BENCH_ITERS, bad, (double)got[0]); + if (ms > 0.0f) { + double flops = 2.0 * (double)BENCH_THREADS * (double)BENCH_ITERS; + + printf("TIMING bench ms=%.3f gflops=%.2f\n", (double)ms, + flops / ((double)ms * 1.0e6)); + } + + cudaEventDestroy(t0); + cudaEventDestroy(t1); + cudaFree(d_out); + free(got); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_churn.cu b/tools/cuda-ab/p2/src/p2_churn.cu new file mode 100644 index 0000000..7d5f9ef --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_churn.cu @@ -0,0 +1,100 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_churn - cudaMalloc/cudaFree churn between launches. + * + * Six iterations allocate 1,2,3,4,1,2 buffers of 2048 ints (8 KB each), + * fill each with a seed-derived kernel, verify, then free all of them. + * A final fresh allocation + launch proves the allocator remains healthy. + * Peak live device memory: 32 KB (4 pages); 14 allocations overall. + */ +#include "common.h" + +#define CH_N 2048 +#define ITERS 6 +#define FINAL 7 + +__global__ void fill(int *p, int seed, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + p[i] = seed * 1000 + i; +} + +int main(int argc, char **argv) +{ + int *got = (int *)malloc(sizeof(int) * CH_N); + int bad = 0; + + if (!got) + return 1; + + if (p2_emit_mode(argc, argv)) { + for (int it = 0; it < ITERS; it++) { + int nbuf = 1 + (it % 4); + + printf("RESULT churn iter=%d nbuf=%d n=%d bad=0\n", it, + nbuf, CH_N); + } + printf("RESULT churn final_buf=1 n=%d bad=0\n", CH_N); + printf("RESULT churn iters=%d allocs=%d bad=0\n", ITERS, + 1 + 2 + 3 + 4 + 1 + 2 + 1); + free(got); + return p2_status(0); + } + + p2_device_info(); + + for (int it = 0; it < ITERS; it++) { + int nbuf = 1 + (it % 4); + int *d[4] = { NULL, NULL, NULL, NULL }; + int sbad = 0, seed = it + 1; + + for (int k = 0; k < nbuf; k++) { + P2_CHECK(cudaMalloc(&d[k], sizeof(int) * CH_N)); + fill<<<(CH_N + 255) / 256, 256>>>(d[k], seed, CH_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d[k], sizeof(int) * CH_N, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < CH_N; i++) + if (got[i] != seed * 1000 + i) + sbad++; + } + for (int k = 0; k < nbuf; k++) + P2_CHECK(cudaFree(d[k])); + printf("RESULT churn iter=%d nbuf=%d n=%d bad=%d\n", it, nbuf, + CH_N, sbad); + bad += sbad; + } + + { + int *d = NULL; + int fbad = 0; + + P2_CHECK(cudaMalloc(&d, sizeof(int) * CH_N)); + fill<<<(CH_N + 255) / 256, 256>>>(d, FINAL, CH_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d, sizeof(int) * CH_N, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < CH_N; i++) + if (got[i] != FINAL * 1000 + i) + fbad++; + P2_CHECK(cudaFree(d)); + printf("RESULT churn final_buf=1 n=%d bad=%d\n", CH_N, fbad); + bad += fbad; + } + + cudaError_t fin = cudaGetLastError(); + + if (fin != cudaSuccess) { + printf("ERR churn final_rc=%d name=%s\n", (int)fin, + cudaGetErrorString(fin)); + bad++; + } + printf("RESULT churn iters=%d allocs=%d bad=%d\n", ITERS, + 1 + 2 + 3 + 4 + 1 + 2 + 1, bad); + free(got); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_managed_probe.cu b/tools/cuda-ab/p2/src/p2_managed_probe.cu new file mode 100644 index 0000000..0d79352 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_managed_probe.cu @@ -0,0 +1,81 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_managed_probe - EXPECTED-FAILURE PROBE (documentation, not acceptance). + * + * Managed (UVM) memory is a known unsupported path for the emulated vGPU. + * This probe attempts cudaMallocManaged, records the outcome as a PROBE + * line, clears any error, and then proves that ordinary cudaMalloc + launch + * still works afterwards. The emulated guest is expected to report a + * failure here; the host oracle may report success. p2/run_guest.sh treats + * this program as informational (content not diffed), but requires the + * program to complete with STATUS OK - i.e. cleanly report the managed + * behaviour and keep working - so a hang, crash or poisoned context still + * fails the suite. + * + * Device memory: 4 KB managed (if it works) + 4 KB normal. + */ +#include "common.h" + +__global__ void fillk(int *p, int v, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + p[i] = v + i; +} + +int main(int argc, char **argv) +{ + if (p2_emit_mode(argc, argv)) { + printf("PROBE managed state=no-oracle-gpu note=emit-mode\n"); + printf("STATUS OK\n"); + return 0; + } + + p2_device_info(); + + int *p = NULL; + cudaError_t e = cudaMallocManaged(&p, sizeof(int) * 4); + + if (e == cudaSuccess) { + cudaError_t e2, e3; + int ok = 0; + + fillk<<<1, 1>>>(p, 5, 4); + e2 = cudaGetLastError(); + e3 = cudaDeviceSynchronize(); + if (e2 == cudaSuccess && e3 == cudaSuccess) + ok = (p[0] == 5 && p[1] == 6 && p[2] == 7 && + p[3] == 8) ? 1 : 0; + printf("PROBE managed alloc=ok write_read=%s\n", + ok ? "ok" : "bad"); + cudaFree(p); + } else { + printf("PROBE managed alloc=fail rc=%d name=%s\n", (int)e, + cudaGetErrorString(e)); + (void)cudaGetLastError(); + } + + int *d = NULL, bad = 0; + cudaError_t e4 = cudaMalloc(&d, sizeof(int) * 4); + + if (e4 == cudaSuccess) { + int h[4]; + + fillk<<<1, 1>>>(d, 100, 4); + (void)cudaGetLastError(); + cudaDeviceSynchronize(); + cudaMemcpy(h, d, sizeof(h), cudaMemcpyDeviceToHost); + for (int i = 0; i < 4; i++) + if (h[i] != 100 + i) + bad++; + cudaFree(d); + printf("PROBE fallback rc=0 bad=%d\n", bad); + } else { + printf("PROBE fallback rc=%d name=%s\n", (int)e4, + cudaGetErrorString(e4)); + bad = 1; + } + + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_mixed.cu b/tools/cuda-ab/p2/src/p2_mixed.cu new file mode 100644 index 0000000..56f993d --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_mixed.cu @@ -0,0 +1,273 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_mixed - mail + vadd + int reduction + 2D grid + 2 streams in one + * process, freeing each phase before the next so the peak live device + * footprint stays at the vadd phase (3 x 32 KB = 24 pages). + * + * phase 1: device-authored mail via a parameter-buffer kernel + * phase 2: vectorAdd n=8192 (element-exact fp32) + * phase 3: 4096 -> 1 int reduction (shared memory + atomicAdd) + * phase 4: 2D grid hash kernel (40x20, block 16x16, partial tiles) + * phase 5: 2 streams, addv + scale per stream (exact fp32) + */ +#include "common.h" + +#define MIX_MAIL "the little gpu man says p2 mixed" +#define VN 8192 +#define RD_N 4096 +#define G2_NX 40 +#define G2_NY 20 +#define ST_N 1024 +#define ST_NS 2 + +__global__ void mail_kernel(char *out, const char *msg, int len) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < len) + out[i] = msg[i]; +} + +__global__ void vadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +__global__ void reduce_sum(const int *in, int *out, int n) +{ + __shared__ int sdata[256]; + int tid = threadIdx.x; + int i = blockIdx.x * blockDim.x + tid; + + sdata[tid] = (i < n) ? in[i] : 0; + __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (tid < s) + sdata[tid] += sdata[tid + s]; + __syncthreads(); + } + if (tid == 0) + atomicAdd(out, sdata[0]); +} + +__global__ void grid2d(int *out, int nx, int ny) +{ + int ix = blockIdx.x * blockDim.x + threadIdx.x; + int iy = blockIdx.y * blockDim.y + threadIdx.y; + + if (ix < nx && iy < ny) { + int idx = iy * nx + ix; + + out[idx] = 1000000 + idx; + } +} + +__global__ void addv(float *x, float v, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + x[i] = x[i] + v; +} + +__global__ void scale(float *x, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + x[i] = x[i] * s; +} + +static int phase_mail(void) +{ + char h[128]; + char *d_out = NULL, *d_msg = NULL; + int len = (int)strlen(MIX_MAIL); + int bad = 0; + + if (cudaMalloc(&d_out, len) != cudaSuccess) + return 1; + if (cudaMalloc(&d_msg, len) != cudaSuccess) { + cudaFree(d_out); + return 1; + } + cudaMemcpy(d_msg, MIX_MAIL, len, cudaMemcpyHostToDevice); + mail_kernel<<<1, 64>>>(d_out, d_msg, len); + cudaGetLastError(); + cudaDeviceSynchronize(); + cudaMemcpy(h, d_out, len, cudaMemcpyDeviceToHost); + bad = (memcmp(h, MIX_MAIL, len) != 0) ? 1 : 0; + if (bad == 0) + printf("MAIL %s\n", h); + cudaFree(d_out); + cudaFree(d_msg); + printf("RESULT mixed mail bad=%d\n", bad); + return bad; +} + +static int phase_vadd(void) +{ + float *h_a = (float *)malloc(sizeof(float) * VN); + float *h_b = (float *)malloc(sizeof(float) * VN); + float *h_c = (float *)malloc(sizeof(float) * VN); + float *d_a = NULL, *d_b = NULL, *d_c = NULL; + int bad = 0; + + if (!h_a || !h_b || !h_c) + return 1; + for (int i = 0; i < VN; i++) { + h_a[i] = (float)(i % 1000) * 0.5f; + h_b[i] = (float)(i % 997) * 0.25f; + } + if (cudaMalloc(&d_a, sizeof(float) * VN) != cudaSuccess) + return 1; + if (cudaMalloc(&d_b, sizeof(float) * VN) != cudaSuccess) + return 1; + if (cudaMalloc(&d_c, sizeof(float) * VN) != cudaSuccess) + return 1; + cudaMemcpy(d_a, h_a, sizeof(float) * VN, cudaMemcpyHostToDevice); + cudaMemcpy(d_b, h_b, sizeof(float) * VN, cudaMemcpyHostToDevice); + vadd<<<(VN + 255) / 256, 256>>>(d_a, d_b, d_c, VN); + cudaGetLastError(); + cudaDeviceSynchronize(); + cudaMemcpy(h_c, d_c, sizeof(float) * VN, cudaMemcpyDeviceToHost); + for (int i = 0; i < VN; i++) + if (h_c[i] != h_a[i] + h_b[i]) + bad++; + cudaFree(d_a); + cudaFree(d_b); + cudaFree(d_c); + free(h_a); + free(h_b); + free(h_c); + printf("RESULT mixed vadd n=%d bad=%d\n", VN, bad); + return bad; +} + +static int phase_reduce(void) +{ + int *h_in = (int *)malloc(sizeof(int) * RD_N); + int *d_in = NULL, *d_out = NULL; + int got = 0, exp = 0, bad = 0; + + if (!h_in) + return 1; + for (int i = 0; i < RD_N; i++) { + h_in[i] = i; + exp += i; + } + if (cudaMalloc(&d_in, sizeof(int) * RD_N) != cudaSuccess) + return 1; + if (cudaMalloc(&d_out, sizeof(int)) != cudaSuccess) + return 1; + cudaMemcpy(d_in, h_in, sizeof(int) * RD_N, cudaMemcpyHostToDevice); + cudaMemset(d_out, 0, sizeof(int)); + reduce_sum<<<(RD_N + 255) / 256, 256>>>(d_in, d_out, RD_N); + cudaGetLastError(); + cudaDeviceSynchronize(); + cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost); + bad = (got != exp) ? 1 : 0; + cudaFree(d_in); + cudaFree(d_out); + free(h_in); + printf("RESULT mixed red n=%d sum=%d bad=%d\n", RD_N, got, bad); + return bad; +} + +static int phase_grid2d(void) +{ + int n = G2_NX * G2_NY; + int *d_out = NULL, *got = (int *)malloc(sizeof(int) * n); + int bad = 0; + + if (!got) + return 1; + if (cudaMalloc(&d_out, sizeof(int) * n) != cudaSuccess) { + free(got); + return 1; + } + grid2d<<>>(d_out, G2_NX, G2_NY); + cudaGetLastError(); + cudaDeviceSynchronize(); + cudaMemcpy(got, d_out, sizeof(int) * n, cudaMemcpyDeviceToHost); + for (int i = 0; i < n; i++) + if (got[i] != 1000000 + i) + bad++; + cudaFree(d_out); + free(got); + printf("RESULT mixed grid2d nx=%d ny=%d bad=%d\n", G2_NX, G2_NY, + bad); + return bad; +} + +static int phase_streams(void) +{ + float *h = (float *)malloc(sizeof(float) * ST_N); + float *got = (float *)malloc(sizeof(float) * ST_N); + cudaStream_t st[ST_NS]; + float *d_x[ST_NS]; + int bad = 0; + + if (!h || !got) + return 1; + for (int i = 0; i < ST_N; i++) + h[i] = (float)(i % 100); + for (int k = 0; k < ST_NS; k++) { + if (cudaStreamCreate(&st[k]) != cudaSuccess) + return 1; + if (cudaMalloc(&d_x[k], sizeof(float) * ST_N) != cudaSuccess) + return 1; + cudaMemcpyAsync(d_x[k], h, sizeof(float) * ST_N, + cudaMemcpyHostToDevice, st[k]); + addv<<<(ST_N + 255) / 256, 256, 0, st[k]>>>(d_x[k], 7.0f, ST_N); + scale<<<(ST_N + 255) / 256, 256, 0, st[k]>>>(d_x[k], 3.0f, ST_N); + cudaGetLastError(); + } + for (int k = 0; k < ST_NS; k++) { + cudaStreamSynchronize(st[k]); + cudaMemcpy(got, d_x[k], sizeof(float) * ST_N, + cudaMemcpyDeviceToHost); + for (int i = 0; i < ST_N; i++) + if (got[i] != ((float)(i % 100) + 7.0f) * 3.0f) + bad++; + cudaStreamDestroy(st[k]); + cudaFree(d_x[k]); + } + free(h); + free(got); + printf("RESULT mixed stream nstreams=%d n=%d bad=%d\n", ST_NS, ST_N, + bad); + return bad; +} + +int main(int argc, char **argv) +{ + if (p2_emit_mode(argc, argv)) { + printf("MAIL %s\n", MIX_MAIL); + printf("RESULT mixed mail bad=0\n"); + printf("RESULT mixed vadd n=%d bad=0\n", VN); + printf("RESULT mixed red n=%d sum=%d bad=0\n", RD_N, + RD_N * (RD_N - 1) / 2); + printf("RESULT mixed grid2d nx=%d ny=%d bad=0\n", G2_NX, G2_NY); + printf("RESULT mixed stream nstreams=%d n=%d bad=0\n", ST_NS, + ST_N); + printf("RESULT mixed total bad=0\n"); + return p2_status(0); + } + + p2_device_info(); + + int bad = 0; + + bad += phase_mail(); + bad += phase_vadd(); + bad += phase_reduce(); + bad += phase_grid2d(); + bad += phase_streams(); + + printf("RESULT mixed total bad=%d\n", bad); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_params.cu b/tools/cuda-ab/p2/src/p2_params.cu new file mode 100644 index 0000000..5dd942c --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_params.cu @@ -0,0 +1,185 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_params - kernel parameter-count and parameter-type coverage. + * + * k0..k12 have 0..12 value parameters respectively (plus the output + * pointer), mixing int, float, double and device-pointer types. Every + * kernel writes one int slot of a 13-int device buffer; the CPU reference + * mirrors each expression exactly and the buffer must match slot for slot. + * + * Device memory: out 52 B, p 8 B, pf 4 B (1 page). + */ +#include "common.h" + +#define NK 13 + +__global__ void k0(int *out) +{ + out[0] = 7; +} + +__global__ void k1(int *out, int a) +{ + out[1] = a + 1; +} + +__global__ void k2(int *out, int a, int b) +{ + out[2] = a + b; +} + +__global__ void k3(int *out, int a, int b, float f1) +{ + out[3] = a + b + (int)f1; +} + +__global__ void k4(int *out, int a, int b, float f1, double c1) +{ + out[4] = a + b + (int)f1 + (int)c1; +} + +__global__ void k5(int *out, int a, int b, float f1, double c1, int d) +{ + out[5] = a + b + (int)f1 + (int)c1 + d; +} + +__global__ void k6(int *out, int a, int b, float f1, double c1, int d, + const int *p) +{ + out[6] = a + b + (int)f1 + (int)c1 + d + p[0]; +} + +__global__ void k7(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q) +{ + out[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0]; +} + +__global__ void k8(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n) +{ + out[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n; +} + +__global__ void k9(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2) +{ + out[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2; +} + +__global__ void k10(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, + double c2) +{ + out[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2; +} + +__global__ void k11(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, + double c2, float s) +{ + out[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s; +} + +__global__ void k12(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, + double c2, float s, int e) +{ + out[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s + e; +} + +/* CPU mirror of the 13 kernel expressions; keeps expected values exact. */ +static void expected(int *exp, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, double c2, + float s, int e) +{ + exp[0] = 7; + exp[1] = a + 1; + exp[2] = a + b; + exp[3] = a + b + (int)f1; + exp[4] = a + b + (int)f1 + (int)c1; + exp[5] = a + b + (int)f1 + (int)c1 + d; + exp[6] = a + b + (int)f1 + (int)c1 + d + p[0]; + exp[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0]; + exp[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n; + exp[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2; + exp[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2; + exp[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s; + exp[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s + e; +} + +static int checksum(const int *v) +{ + int sum = 0; + + for (int i = 0; i < NK; i++) + sum += v[i]; + return sum; +} + +int main(int argc, char **argv) +{ + int a = 3, b = 5, d = 11, n = 2, e = 4; + float f1 = 6.5f, f2 = 8.5f, s = 2.5f; + double c1 = 2.5, c2 = 3.5; + int hp[2] = { 5, 6 }, hq[2] = { 6, 7 }; + int exp[NK]; + + expected(exp, a, b, f1, c1, d, hp, hq, n, f2, c2, s, e); + + if (p2_emit_mode(argc, argv)) { + printf("RESULT params kernels=%d nparam_max=12 bad=0 checksum=%d\n", + NK, checksum(exp)); + return p2_status(0); + } + + p2_device_info(); + + int *d_out = NULL, *got = (int *)malloc(sizeof(int) * NK); + int *dp = NULL, *dq = NULL; + + if (!got) + return 1; + P2_CHECK(cudaMalloc(&d_out, sizeof(int) * NK)); + P2_CHECK(cudaMalloc(&dp, sizeof(int) * 2)); + P2_CHECK(cudaMalloc(&dq, sizeof(int) * 2)); + P2_CHECK(cudaMemcpy(dp, hp, sizeof(int) * 2, cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemcpy(dq, hq, sizeof(int) * 2, cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemset(d_out, 0, sizeof(int) * NK)); + + k0<<<1, 1>>>(d_out); + k1<<<1, 1>>>(d_out, a); + k2<<<1, 1>>>(d_out, a, b); + k3<<<1, 1>>>(d_out, a, b, f1); + k4<<<1, 1>>>(d_out, a, b, f1, c1); + k5<<<1, 1>>>(d_out, a, b, f1, c1, d); + k6<<<1, 1>>>(d_out, a, b, f1, c1, d, dp); + k7<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq); + k8<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n); + k9<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2); + k10<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2); + k11<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s); + k12<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s, e); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(int) * NK, + cudaMemcpyDeviceToHost)); + + int bad = 0; + + for (int i = 0; i < NK; i++) + if (got[i] != exp[i]) + bad++; + printf("RESULT params kernels=%d nparam_max=12 bad=%d checksum=%d\n", + NK, bad, checksum(got)); + + cudaFree(d_out); cudaFree(dp); cudaFree(dq); + free(got); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_printf_err.cu b/tools/cuda-ab/p2/src/p2_printf_err.cu new file mode 100644 index 0000000..d7f449e --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_printf_err.cu @@ -0,0 +1,100 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_printf_err - device printf plus an intentional launch error path. + * + * 1. A 4-thread kernel prints deterministic P2PRINT lines from the device. + * 2. A kernel launched with gridDim.x == 0 must be rejected: + * cudaErrorInvalidConfiguration (9), and the error must be cleared by + * cudaGetLastError. A blockDim.x == 0 launch is exercised too. + * 3. A normal launch after the errors proves error recovery. + * + * Expected error strings come from the statically linked CUDA runtime, so + * they must be identical on host and guest. + * Device memory: out 4 B. + */ +#include "common.h" + +__global__ void hello(int nthreads) +{ + int i = threadIdx.x; + + if (i < nthreads) + printf("P2PRINT hello i=%d\n", i); +} + +__global__ void second(void) +{ + if (threadIdx.x == 0) + printf("P2PRINT second c=%c v=%d\n", 'Z', 7); +} + +__global__ void simple(int *out) +{ + out[0] = 42; +} + +int main(int argc, char **argv) +{ + if (p2_emit_mode(argc, argv)) { + printf("P2PRINT hello i=0\n"); + printf("P2PRINT hello i=1\n"); + printf("P2PRINT hello i=2\n"); + printf("P2PRINT hello i=3\n"); + printf("P2PRINT second c=Z v=7\n"); + printf("ERR launch grid0 rc=%d name=%s\n", + (int)cudaErrorInvalidConfiguration, + cudaGetErrorString(cudaErrorInvalidConfiguration)); + printf("ERR launch block0 rejected=1\n"); + printf("ERR cleared rc=0\n"); + printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n"); + return p2_status(0); + } + + p2_device_info(); + + int *d_out = NULL; + P2_CHECK(cudaMalloc(&d_out, sizeof(int))); + + hello<<<1, 4>>>(4); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + second<<<1, 4>>>(); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + + (void)cudaGetLastError(); + + int badgrid = 0, badblock = 0; + cudaError_t e_grid, e_block, cleared; + int got = 0; + + simple<<>>(d_out); + e_grid = cudaGetLastError(); + printf("ERR launch grid0 rc=%d name=%s\n", (int)e_grid, + cudaGetErrorString(e_grid)); + + simple<<<1, badblock>>>(d_out); + e_block = cudaGetLastError(); + printf("ERR launch block0 rejected=%d\n", + (e_block != cudaSuccess) ? 1 : 0); + + cleared = cudaGetLastError(); + printf("ERR cleared rc=%d\n", (int)cleared); + + simple<<<1, 1>>>(d_out); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost)); + + int grid0_bad = (e_grid != cudaErrorInvalidConfiguration) ? 1 : 0; + int block0_bad = (e_block == cudaSuccess) ? 1 : 0; + int clear_bad = (cleared != cudaSuccess) ? 1 : 0; + int recovered = (got == 42) ? 1 : 0; + int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1); + + printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n", + grid0_bad, block0_bad, clear_bad, recovered); + + cudaFree(d_out); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_seq5.cu b/tools/cuda-ab/p2/src/p2_seq5.cu new file mode 100644 index 0000000..8a778f4 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_seq5.cu @@ -0,0 +1,127 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_seq5 - five sequential kernel launches in one process on the same + * buffers, verifying the device state after every launch. + * + * step 1: c = a + b + * step 2: c = c + a + * step 3: c = c + b + * step 4: c = c * 2 + * step 5: c = c + 1 + * + * Inputs are small integers so every operation is exact in fp32; the CPU + * reference reproduces the same chain and must match element for element. + * Device memory: 3 x 4 KB pointers (3 pages total). + */ +#include "common.h" + +#define N 1024 + +__global__ void kadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +__global__ void kmul2(float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * 2.0f; +} + +__global__ void kaddv(float *c, float v, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] + v; +} + +static void make_inputs(float *a, float *b) +{ + for (int i = 0; i < N; i++) { + a[i] = (float)(i % 7); + b[i] = (float)(i % 5); + } +} + +static void step_cpu(float *c, const float *a, const float *b, int step) +{ + for (int i = 0; i < N; i++) { + switch (step) { + case 1: c[i] = a[i] + b[i]; break; + case 2: c[i] = c[i] + a[i]; break; + case 3: c[i] = c[i] + b[i]; break; + case 4: c[i] = c[i] * 2.0f; break; + case 5: c[i] = c[i] + 1.0f; break; + } + } +} + +int main(int argc, char **argv) +{ + float *h_a, *h_b, *h_c, *h_exp; + int bad = 0; + + if (p2_emit_mode(argc, argv)) { + h_a = (float *)malloc(sizeof(float) * N); + h_b = (float *)malloc(sizeof(float) * N); + h_exp = (float *)malloc(sizeof(float) * N); + make_inputs(h_a, h_b); + for (int step = 1; step <= 5; step++) { + step_cpu(h_exp, h_a, h_b, step); + printf("RESULT seq step=%d n=%d bad=0\n", step, N); + } + printf("RESULT seq launches=5 n=%d bad=0\n", N); + free(h_a); free(h_b); free(h_exp); + return p2_status(0); + } + + p2_device_info(); + h_a = (float *)malloc(sizeof(float) * N); + h_b = (float *)malloc(sizeof(float) * N); + h_c = (float *)malloc(sizeof(float) * N); + h_exp = (float *)malloc(sizeof(float) * N); + if (!h_a || !h_b || !h_c || !h_exp) + return 1; + make_inputs(h_a, h_b); + + float *d_a = NULL, *d_b = NULL, *d_c = NULL; + + P2_CHECK(cudaMalloc(&d_a, sizeof(float) * N)); + P2_CHECK(cudaMalloc(&d_b, sizeof(float) * N)); + P2_CHECK(cudaMalloc(&d_c, sizeof(float) * N)); + P2_CHECK(cudaMemcpy(d_a, h_a, sizeof(float) * N, cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemcpy(d_b, h_b, sizeof(float) * N, cudaMemcpyHostToDevice)); + + for (int step = 1; step <= 5; step++) { + int sbad = 0; + + switch (step) { + case 1: kadd<<<(N + 255) / 256, 256>>>(d_a, d_b, d_c, N); break; + case 2: kadd<<<(N + 255) / 256, 256>>>(d_c, d_a, d_c, N); break; + case 3: kadd<<<(N + 255) / 256, 256>>>(d_c, d_b, d_c, N); break; + case 4: kmul2<<<(N + 255) / 256, 256>>>(d_c, N); break; + case 5: kaddv<<<(N + 255) / 256, 256>>>(d_c, 1.0f, N); break; + } + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(h_c, d_c, sizeof(float) * N, + cudaMemcpyDeviceToHost)); + step_cpu(h_exp, h_a, h_b, step); + for (int i = 0; i < N; i++) + if (h_c[i] != h_exp[i]) + sbad++; + printf("RESULT seq step=%d n=%d bad=%d\n", step, N, sbad); + bad += sbad; + } + printf("RESULT seq launches=5 n=%d bad=%d\n", N, bad); + + cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); + free(h_a); free(h_b); free(h_c); free(h_exp); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_shapes.cu b/tools/cuda-ab/p2/src/p2_shapes.cu new file mode 100644 index 0000000..93ae0e9 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_shapes.cu @@ -0,0 +1,111 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_shapes - grid/block shape variation with a bound-checked index-hash + * kernel: + * case 1: 1x1x1 grid, 1x1x1 block (minimum launch) + * case 2: odd sizes (n=1013, block 33, 31 blocks) + * case 3: non-multiple n (n=1000, block 64, partial last block) + * case 4: 2D grid (40x20, block 16x16, partial tiles) + * case 5: 3D grid (24x8x8, block 8x4x2) + * + * Device memory: one 8 KB int buffer reused across cases (2 pages). + */ +#include "common.h" + +#define MAXN 2048 + +struct shape { + int nx, ny, nz; /* logical extent */ + int bx, by, bz; /* block dims */ + int gx, gy, gz; /* grid dims */ + int tag; +}; + +static const struct shape cases[] = { + { 1, 1, 1, 1, 1, 1, 1, 1, 1, 1 }, + {1013, 1, 1, 33, 1, 1, 31, 1, 1, 2 }, + {1000, 1, 1, 64, 1, 1, 16, 1, 1, 3 }, + { 40, 20, 1, 16, 16, 1, 5, 2, 1, 4 }, + { 24, 8, 8, 8, 4, 2, 3, 2, 4, 5 }, +}; + +#define NCASES ((int)(sizeof(cases) / sizeof(cases[0]))) + +__global__ void shape_kernel(int *out, int nx, int ny, int nz, int tag) +{ + int ix = blockIdx.x * blockDim.x + threadIdx.x; + int iy = blockIdx.y * blockDim.y + threadIdx.y; + int iz = blockIdx.z * blockDim.z + threadIdx.z; + + if (ix < nx && iy < ny && iz < nz) { + int idx = (iz * ny + iy) * nx + ix; + + out[idx] = tag * 1000000 + idx; + } +} + +static void shape_expected(int *exp, const struct shape *s) +{ + for (int iz = 0; iz < s->nz; iz++) + for (int iy = 0; iy < s->ny; iy++) + for (int ix = 0; ix < s->nx; ix++) { + int idx = (iz * s->ny + iy) * s->nx + ix; + + exp[idx] = s->tag * 1000000 + idx; + } +} + +int main(int argc, char **argv) +{ + int bad = 0; + + if (p2_emit_mode(argc, argv)) { + int *exp = (int *)malloc(sizeof(int) * MAXN); + + for (int c = 0; c < NCASES; c++) { + int n = cases[c].nx * cases[c].ny * cases[c].nz; + + shape_expected(exp, &cases[c]); + printf("RESULT shapes case=%d n=%d bad=0\n", + cases[c].tag, n); + } + printf("RESULT shapes cases=%d bad=0\n", NCASES); + free(exp); + return p2_status(0); + } + + p2_device_info(); + + int *d_out = NULL, *h_got = (int *)malloc(sizeof(int) * MAXN); + int *h_exp = (int *)malloc(sizeof(int) * MAXN); + + if (!h_got || !h_exp) + return 1; + P2_CHECK(cudaMalloc(&d_out, sizeof(int) * MAXN)); + + for (int c = 0; c < NCASES; c++) { + const struct shape *s = &cases[c]; + int n = s->nx * s->ny * s->nz; + dim3 grid(s->gx, s->gy, s->gz), block(s->bx, s->by, s->bz); + int sbad = 0; + + shape_kernel<<>>(d_out, s->nx, s->ny, s->nz, + s->tag); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(h_got, d_out, sizeof(int) * n, + cudaMemcpyDeviceToHost)); + memset(h_exp, 0, sizeof(int) * n); + shape_expected(h_exp, s); + for (int i = 0; i < n; i++) + if (h_got[i] != h_exp[i]) + sbad++; + printf("RESULT shapes case=%d n=%d bad=%d\n", s->tag, n, sbad); + bad += sbad; + } + printf("RESULT shapes cases=%d bad=%d\n", NCASES, bad); + + cudaFree(d_out); + free(h_got); free(h_exp); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_smem.cu b/tools/cuda-ab/p2/src/p2_smem.cu new file mode 100644 index 0000000..5db0d52 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_smem.cu @@ -0,0 +1,144 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_smem - shared-memory kernels: + * reduce_sum: 4096 int32 -> single sum, 16 blocks x 256 threads, shared + * tree reduction plus a global atomicAdd. + * gemm_tiled: 32x32 int32 GEMM with 16x16 shared-memory tiles, + * grid(2,2) x block(16,16). + * + * All arithmetic is exact in int32 (inputs are small non-negative ints), so + * device results must equal the CPU reference bit for bit. + * Device memory: in 16 KB, out 4 B, then A/B/C 4 KB each (peak 5 pages). + */ +#include "common.h" + +#define R_N 4096 +#define TS 16 +#define M 32 +#define K 32 +#define NN 32 + +__global__ void reduce_sum(const int *in, int *out, int n) +{ + __shared__ int sdata[256]; + int tid = threadIdx.x; + int i = blockIdx.x * blockDim.x + tid; + + sdata[tid] = (i < n) ? in[i] : 0; + __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (tid < s) + sdata[tid] += sdata[tid + s]; + __syncthreads(); + } + if (tid == 0) + atomicAdd(out, sdata[0]); +} + +__global__ void gemm_tiled(const int *A, const int *B, int *C) +{ + __shared__ int As[TS][TS], Bs[TS][TS]; + int row = blockIdx.y * TS + threadIdx.y; + int col = blockIdx.x * TS + threadIdx.x; + int sum = 0; + + for (int t = 0; t < K / TS; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TS + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = B[(t * TS + threadIdx.y) * NN + col]; + __syncthreads(); + for (int j = 0; j < TS; j++) + sum += As[threadIdx.y][j] * Bs[j][threadIdx.x]; + __syncthreads(); + } + C[row * NN + col] = sum; +} + +int main(int argc, char **argv) +{ + int *h_in = (int *)malloc(sizeof(int) * R_N); + int *h_A = (int *)malloc(sizeof(int) * M * K); + int *h_B = (int *)malloc(sizeof(int) * K * NN); + int *h_C = (int *)malloc(sizeof(int) * M * NN); + int red_sum = 0, gemm_sum = 0, bad = 0; + + if (!h_in || !h_A || !h_B || !h_C) + return 1; + for (int i = 0; i < R_N; i++) + h_in[i] = i; + for (int i = 0; i < M * K; i++) + h_A[i] = i % 5; + for (int i = 0; i < K * NN; i++) + h_B[i] = i % 3; + for (int i = 0; i < R_N; i++) + red_sum += h_in[i]; + for (int r = 0; r < M; r++) + for (int c = 0; c < NN; c++) { + int s = 0; + + for (int t = 0; t < K; t++) + s += h_A[r * K + t] * h_B[t * NN + c]; + h_C[r * NN + c] = s; + gemm_sum += s; + } + + if (p2_emit_mode(argc, argv)) { + printf("RESULT smem reduce n=%d sum=%d bad=0\n", R_N, red_sum); + printf("RESULT smem gemm m=%d n=%d k=%d checksum=%d bad=0\n", + M, NN, K, gemm_sum); + free(h_in); free(h_A); free(h_B); free(h_C); + return p2_status(0); + } + + p2_device_info(); + + int *d_in = NULL, *d_out = NULL, *d_A = NULL, *d_B = NULL, *d_C = NULL; + int got_sum = 0, got_gemm = 0; + + P2_CHECK(cudaMalloc(&d_in, sizeof(int) * R_N)); + P2_CHECK(cudaMalloc(&d_out, sizeof(int))); + P2_CHECK(cudaMemcpy(d_in, h_in, sizeof(int) * R_N, + cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemset(d_out, 0, sizeof(int))); + reduce_sum<<<(R_N + 255) / 256, 256>>>(d_in, d_out, R_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(&got_sum, d_out, sizeof(int), + cudaMemcpyDeviceToHost)); + printf("RESULT smem reduce n=%d sum=%d bad=%d\n", R_N, got_sum, + got_sum == red_sum ? 0 : 1); + bad += (got_sum == red_sum) ? 0 : 1; + cudaFree(d_in); cudaFree(d_out); + + P2_CHECK(cudaMalloc(&d_A, sizeof(int) * M * K)); + P2_CHECK(cudaMalloc(&d_B, sizeof(int) * K * NN)); + P2_CHECK(cudaMalloc(&d_C, sizeof(int) * M * NN)); + P2_CHECK(cudaMemcpy(d_A, h_A, sizeof(int) * M * K, + cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemcpy(d_B, h_B, sizeof(int) * K * NN, + cudaMemcpyHostToDevice)); + gemm_tiled<<>>(d_A, d_B, d_C); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + + int *got_C = (int *)malloc(sizeof(int) * M * NN); + + if (!got_C) + return 1; + P2_CHECK(cudaMemcpy(got_C, d_C, sizeof(int) * M * NN, + cudaMemcpyDeviceToHost)); + int cbad = 0; + + for (int i = 0; i < M * NN; i++) { + got_gemm += got_C[i]; + if (got_C[i] != h_C[i]) + cbad++; + } + printf("RESULT smem gemm m=%d n=%d k=%d checksum=%d bad=%d\n", + M, NN, K, got_gemm, cbad); + bad += cbad; + + cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); + free(got_C); + free(h_in); free(h_A); free(h_B); free(h_C); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_streams.cu b/tools/cuda-ab/p2/src/p2_streams.cu new file mode 100644 index 0000000..09d6ed3 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_streams.cu @@ -0,0 +1,109 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_streams - three created streams, two kernels per stream (scale then + * add), independent 1024-float buffers, per-stream events. + * + * Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32). + * After all streams complete, each event must query as complete. + * Device memory: 3 x 4 KB (3 pages). + */ +#include "common.h" + +#define SN 1024 +#define NS 3 + +__global__ void scale(float *x, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + x[i] = x[i] * s; +} + +__global__ void addv(float *x, float v, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + x[i] = x[i] + v; +} + +static void fill(float *h, int n) +{ + for (int i = 0; i < n; i++) + h[i] = (float)(i % 100); +} + +static float expected1(int i) +{ + return ((float)(i % 100) * 2.0f) + 1.0f; +} + +int main(int argc, char **argv) +{ + float *h = (float *)malloc(sizeof(float) * SN); + float *got = (float *)malloc(sizeof(float) * SN); + int bad = 0; + + if (!h || !got) + return 1; + fill(h, SN); + + if (p2_emit_mode(argc, argv)) { + printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n", + NS, NS * 2, SN); + printf("RESULT streams events=%d bad=0\n", NS); + free(h); free(got); + return p2_status(0); + } + + p2_device_info(); + + cudaStream_t st[NS]; + cudaEvent_t ev[NS]; + float *d_x[NS]; + + for (int k = 0; k < NS; k++) { + P2_CHECK(cudaStreamCreate(&st[k])); + P2_CHECK(cudaEventCreate(&ev[k])); + P2_CHECK(cudaMalloc(&d_x[k], sizeof(float) * SN)); + P2_CHECK(cudaMemcpyAsync(d_x[k], h, sizeof(float) * SN, + cudaMemcpyHostToDevice, st[k])); + scale<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 2.0f, SN); + addv<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 1.0f, SN); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaEventRecord(ev[k], st[k])); + } + + int ibad = 0; + + for (int k = 0; k < NS; k++) { + cudaError_t q; + int sbad = 0; + + P2_CHECK(cudaStreamSynchronize(st[k])); + P2_CHECK(cudaMemcpy(got, d_x[k], sizeof(float) * SN, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < SN; i++) + if (got[i] != expected1(i)) + sbad++; + q = cudaEventQuery(ev[k]); + if (q != cudaSuccess) + sbad++; + printf("RESULT streams stream=%d query_rc=%d bad=%d\n", k, (int)q, + sbad); + ibad += sbad; + } + printf("RESULT streams nstreams=%d launches=%d n=%d bad=%d\n", + NS, NS * 2, SN, ibad); + printf("RESULT streams events=%d bad=%d\n", NS, ibad == 0 ? 0 : 1); + bad = ibad; + + for (int k = 0; k < NS; k++) { + cudaEventDestroy(ev[k]); + cudaStreamDestroy(st[k]); + cudaFree(d_x[k]); + } + free(h); free(got); + return p2_status(bad); +} diff --git a/tools/guest-diag/guest_pin_stream.py b/tools/guest-diag/guest_pin_stream.py index f0f07d8..581dc9d 100644 --- a/tools/guest-diag/guest_pin_stream.py +++ b/tools/guest-diag/guest_pin_stream.py @@ -13,6 +13,7 @@ import os import re import socket +import struct import subprocess import sys import threading @@ -66,6 +67,371 @@ TRACE_PIPE = ("/sys/kernel/tracing/trace_pipe" SEM_EVERY = int(os.environ.get("STREAM_SEM_EVERY", "4")) +class Kcore: + """Kernel-truth reader: /proc/kcore for contents + manual page-table walk. + + The host's QMP-based struct walk intermittently reads garbage while the + client spins in user mode (PTI hides the kernel half from the sampled + vCPU), which left the CE push-buffer / semaphore-pool pages unpinned on + ~half of all boots (the deterministic cudaMalloc stall). The guest can + read its own kernel memory reliably: /proc/kcore has an ELF segment for + every kernel mapping, and the direct-map segment lets us walk the page + tables (init_top_pgt -> ... -> PTE) to turn the vmalloc CPU alias of a + UVM rm_mem into the guest physical addresses the host needs. + """ + + def __init__(self): + self.ok = False + self.fd = None + self.segs = [] + self.phys_base = 0 + self.page_offset_base = 0 + self.pgd_pa = 0 + self._pa_cache = {} + + def load(self): + try: + self.fd = open("/proc/kcore", "rb") + hdr = self.fd.read(64) + if hdr[:4] != b"\x7fELF": + return False + e_phoff = struct.unpack_from("> shift) & 0x1FF) * 8) + if e is None or not (e & 1): + self._pa_cache[page] = 0 + return None + if level == 3: + pa = e & 0x000FFFFFFFF000 + break + if e & 0x80: + if level == 1: + pa = (e & 0x000FFFFFC0000000) + (va & ((1 << 30) - 1)) + break + if level == 2: + pa = (e & 0x000FFFFFFFE00000) + (va & ((1 << 21) - 1)) + break + self._pa_cache[page] = 0 + return None + pa = e & 0x000FFFFFFFF000 + self._pa_cache[page] = pa + return pa + + +def is_kern_ptr(v): + return v and (v >> 48) == 0xFFFF + + +def is_gpu_va(v): + return v and 0x120000000 <= v < 0x200000000 + + +def _direct_map_pa(kc, va): + """Direct-map aliases translate linearly: pa = va - page_offset_base.""" + if kc.page_offset_base <= va < kc.page_offset_base + (1 << 40): + return va - kc.page_offset_base + return None + + +def _vmalloc_pages_for(kc, va): + """Page count of the vmap/vmalloc area containing va (0 if unknown).""" + now = time.time() + if not hasattr(_vmalloc_pages_for, "areas") or \ + now - _vmalloc_pages_for.ts > 5.0: + areas = [] + try: + for line in open("/proc/vmallocinfo"): + p = line.split() + if len(p) < 3: + continue + try: + lo, hi = p[0].split("-") + areas.append((int(lo, 16), int(hi, 16))) + except ValueError: + continue + except OSError: + pass + _vmalloc_pages_for.areas = areas + _vmalloc_pages_for.ts = now + base = va & ~0xFFF + for lo, hi in _vmalloc_pages_for.areas: + if lo <= base < hi: + return (hi - base) >> 12 + return 0 + + +def resolve_mem_pins(kc, mem, out, default_pages=16): + """uvm_rm_mem kptr -> GPU-VA/GPA runs (direct map or page-table walk).""" + if not is_kern_ptr(mem): + return 0 + vas = [kc.read64(mem + 0x10 + i * 8) for i in range(16)] + cpu = next((v for v in vas if is_kern_ptr(v)), None) + gpus = [v for v in vas if is_gpu_va(v)] + if not cpu or not gpus: + return 0 + size = kc.read64(mem + 0x60) or 0 + dm = _direct_map_pa(kc, cpu) + emitted = 0 + if dm is not None: + npages = max(1, min(size >> 12, 4098)) if size else default_pages + for gv in gpus: + out.append("0x%x,0x%x,%d\n" % (gv, dm, npages)) + emitted += 1 + return emitted + npages = (max(1, min(size >> 12, 4098)) if size + else _vmalloc_pages_for(kc, cpu)) + if not npages: + npages = 1 + runs = [] + cur = None + for k in range(npages): + pa = kc.va_to_pa(cpu + k * 0x1000) + if pa is None: + if cur: + runs.append(cur) + cur = None + continue + if cur and pa == cur[1] + cur[2] * 0x1000: + cur[2] += 1 + else: + if cur: + runs.append(cur) + cur = [k, pa, 1] + if cur: + runs.append(cur) + for gv in gpus: + for k, pa, n in runs: + out.append("0x%x,0x%x,%d\n" % (gv + k * 0x1000, pa, n)) + emitted += 1 + return emitted + + +def _emit_ring_pins(kc, gpu_va, entries_cpu, entries, out): + """Pin the GPFIFO ring pages (and the pages after it, where the GPPUT + buffer usually sits) from the ring's CPU mapping.""" + emitted = 0 + if not is_kern_ptr(entries_cpu) or not (32 <= entries <= 0x100000): + return 0 + npages = min(((entries * 8) + 0xFFF) // 0x1000 + 2, 8) + runs = [] + cur = None + for k in range(npages): + pa = kc.va_to_pa(entries_cpu + k * 0x1000) + if pa is None: + if cur: + runs.append(cur) + cur = None + continue + if cur and pa == cur[1] + cur[2] * 0x1000: + cur[2] += 1 + else: + if cur: + runs.append(cur) + cur = [k, pa, 1] + if cur: + runs.append(cur) + for k, pa, n in runs: + out.append("0x%x,0x%x,%d\n" % (gpu_va + k * 0x1000, pa, n)) + emitted += 1 + return emitted + + +def resolve_channel_from_sem(kc, sem, out): + """Tracking sem (embedded in uvm_channel_t) -> ring/GPPUT pins. + + uvch events do not fire on this driver build, but the channel's tracking + semaphore is a uvm_channel_t member updated on every push. Locate + channel_info by scanning around the sem for its gpFifoGpuVa value (a + page-aligned GPU VA in the GR ring window), then read the ring/GPPUT + CPU pointers (UvmGpuChannelInfo offsets: gpPut +0x08, gpFifoEntries + +0x10, numGpFifoEntries +0x18, gpFifoGpuVa +0x58, gpPutGpuVa +0x60). + """ + if not is_kern_ptr(sem): + return 0 + lo = sem - 0x2000 + data = kc.read(lo, 0x4000) + if not data: + return 0 + for off in range(0, len(data) - 8, 8): + v = struct.unpack_from(" pin the GPFIFO ring + GPPUT buffers. + + The ring and GP_PUT buffers are CPU-mapped in the guest kernel (UVM + writes pushes and the put counter directly), so their guest-physical + backing is reachable with the page-table walk. Without these pins the + host cannot read GR channel entries (the stall class: RM event loop + waits for a completion whose ring entry was never processed). + """ + if not is_kern_ptr(ch): + return 0 + base = None + gpu_va = None + data = kc.read(ch, 0x800) + if not data: + return 0 + for off in range(0, len(data) - 8, 8): + v = struct.unpack_from(" 0x100000: + return 0 + emitted = 0 + entries_cpu = kc.read64(base + 0x10) + if is_kern_ptr(entries_cpu): + nbytes = entries * 8 + npages = min((nbytes + 0xFFF) // 0x1000 + 2, 8) + runs = [] + cur = None + for k in range(npages): + pa = kc.va_to_pa(entries_cpu + k * 0x1000) + if pa is None: + if cur: + runs.append(cur) + cur = None + continue + if cur and pa == cur[1] + cur[2] * 0x1000: + cur[2] += 1 + else: + if cur: + runs.append(cur) + cur = [k, pa, 1] + if cur: + runs.append(cur) + for k, pa, n in runs: + out.append("0x%x,0x%x,%d\n" % (gpu_va + k * 0x1000, pa, n)) + emitted += 1 + gp_put = kc.read64(base + 0x08) + gp_put_va = kc.read64(base + 0x60) + if is_kern_ptr(gp_put) and gp_put_va and \ + GR_RING_LO <= gp_put_va < 0x201000000: + pa = kc.va_to_pa(gp_put) + if pa is not None: + out.append("0x%x,0x%x,1\n" % (gp_put_va & ~0xFFF, pa & ~0xFFF)) + emitted += 1 + gp_get = kc.read64(base + 0x00) + gp_get_va = kc.read64(base + 0x68) + if is_kern_ptr(gp_get) and gp_get_va and \ + GR_RING_LO <= gp_get_va < 0x201000000: + pa = kc.va_to_pa(gp_get) + if pa is not None: + out.append("0x%x,0x%x,1\n" % (gp_get_va & ~0xFFF, pa & ~0xFFF)) + emitted += 1 + return emitted + + + def clear_trace(): """Drop stale ring events so the fallback tail reader cannot re-learn a previous client's freed sems after a client switch. The new client's @@ -133,6 +499,11 @@ class TraceReader: if ("p", v) not in self.seen: self.seen.add(("p", v)) self.lines.append("#pb 0x%x\n" % v) + for m in re.finditer(rb"ch=0x([0-9a-fA-F]+)", data): + v = int(m.group(1), 16) + if ("c", v) not in self.seen: + self.seen.add(("c", v)) + self.lines.append("#ch 0x%x\n" % v) def drain(self): with self.lock: @@ -163,17 +534,27 @@ def new_sems(seen): def cdt_pid(): prefixes = tuple(p for p in os.environ.get("CLIENT_PREFIXES", "cdt,va,victory").split(",") if p) + best = None for d in os.listdir('/proc'): if not d.isdigit(): continue try: with open('/proc/%s/comm' % d) as f: name = f.read().strip() - if any(name.startswith(p) for p in prefixes): - return int(d) + if not any(name.startswith(p) for p in prefixes): + continue + # Skip zombies: a defunct previous client keeps its comm and + # would otherwise shadow the live client (pagemap reads then fail + # and no pins are ever streamed). + with open('/proc/%s/stat' % d) as f: + st = f.read() + rp = st.rfind(')') + if rp >= 0 and st[rp + 2:rp + 3] == 'Z': + continue + best = int(d) except OSError: pass - return None + return best def scan_window(fd, lo, hi, span, out): @@ -318,7 +699,21 @@ def main(): sock = None seen = {} seen_sems = set() + seen_pbs = set() + seen_chs = set() + resolved_sems = set() + resolved_pbs = set() + resolved_chs = set() + retry_at = {} tracer = TraceReader() + kc = Kcore() + kc_ok = kc.load() + try: + with open('/tmp/guest_pin_stream.log', 'a') as f: + f.write('%s resolver: kcore=%s (sem->pool/pushbuffer pins)\n' + % (time.strftime('%T'), 'ON' if kc_ok else 'OFF')) + except OSError: + pass last_pid = None last_log = 0.0 loop = 0 @@ -329,12 +724,20 @@ def main(): last_pid = None seen = {} seen_sems = set() + seen_pbs = set() + seen_chs = set() time.sleep(INTERVAL) continue if pid != last_pid: last_pid = pid seen = {} seen_sems = set() + seen_pbs = set() + seen_chs = set() + resolved_sems = set() + resolved_pbs = set() + resolved_chs = set() + retry_at = {} clear_trace() if sock is not None: try: @@ -373,9 +776,70 @@ def main(): pending_cur = cur lines = [] for ctl in tracer.drain(): - if ctl.startswith("#sem"): - seen_sems.add(int(ctl.split()[1], 16)) + parts = ctl.split() + if ctl.startswith("#sem") and len(parts) >= 2: + seen_sems.add(int(parts[1], 16)) + elif ctl.startswith("#pb") and len(parts) >= 2: + seen_pbs.add(int(parts[1], 16)) + elif ctl.startswith("#ch") and len(parts) >= 2: + seen_chs.add(int(parts[1], 16)) lines.append(ctl) + if kc_ok: + now_r = time.time() + added = 0 + for sem in list(seen_sems): + if sem in resolved_sems or now_r < retry_at.get(sem, 0): + continue + got = [] + try: + resolve_sem_pins(kc, sem, got) + except Exception: + got = [] + if got: + lines.extend(got) + resolved_sems.add(sem) + added += len(got) + else: + retry_at[sem] = now_r + 1.0 + for pb in list(seen_pbs): + if pb in resolved_pbs or now_r < retry_at.get(pb, 0): + continue + got = [] + try: + resolve_pb_pins(kc, pb, got) + except Exception: + got = [] + if got: + lines.extend(got) + resolved_pbs.add(pb) + added += len(got) + else: + retry_at[pb] = now_r + 1.0 + for ch in list(seen_chs): + if ch in resolved_chs or now_r < retry_at.get(ch, 0): + continue + got = [] + try: + resolve_channel_pins(kc, ch, got) + except Exception: + got = [] + if got: + lines.extend(got) + resolved_chs.add(ch) + added += len(got) + else: + retry_at[ch] = now_r + 1.0 + if added: + try: + with open('/tmp/guest_pin_stream.log', 'a') as f: + f.write('%s resolver pins: +%d run(s) ' + '(sems=%d/%d pbs=%d/%d chs=%d/%d)\n' + % (time.strftime('%T'), added, + len(resolved_sems), len(seen_sems), + len(resolved_pbs), len(seen_pbs), + len(resolved_chs), len(seen_chs))) + except OSError: + pass now = time.time() if now - last_log > 1.0: last_log = now -- cgit v1.2.3