kernel: execute pushed CE work (PTE memsets, copies) and track semaphores

The progress service now decodes push buffers as method streams and performs
the work the driver pushes, not just the tracking semaphore:

- C5B5 method decoder: SET_REMAP_CONST_A/B + COMPONENTS + OFFSET_OUT +
  LINE_LENGTH_IN + LAUNCH_DMA = memset (into vidmem or guest RAM);
  OFFSET_IN + OFFSET_OUT + LAUNCH_DMA = copy; SET_SEMAPHORE_A/B/PAYLOAD +
  LAUNCH_DMA = release. Contiguous method runs (count 2/4) are applied in
  full - OFFSET_IN/OUT arrive as one run.
- Pushes span page boundaries, so pages are decoded as method streams
  without a start marker. Collected operations execute only when the
  push's release is reached and still pending (payload ahead of the
  semaphore), so recycled buffers replaying old pushes are skipped and
  their stale memsets cannot overwrite newer state.
- VA->PA map is first-wins: stale copies of freed allocations and
  unrelated adjacent u64 pairs produced bogus entries that hijacked a
  good mapping and sent semaphore writes to the wrong page.
- dump_guest module parameter (hex dump of guest PA, kthread via
  completion; kthread_stop before the thread runs loses the work).

With this the guest no longer stalls in uvm_tracker_wait and no longer
oopses in the UVM path; the failure moved to a userspace segfault inside
libcuda (NULL-0x38) after the driver's page allocations.

Also add tools/cuda-ab: run one deterministic CUDA workload on the host
(bare metal oracle) and in the guest, then compare the stable result
lines. The binary links the CUDA runtime statically so the guest only
needs the driver's libcuda.so.

5 files changed, 216 insertions, 0 deletions

diff --git a/tools/cuda-ab/README.md b/tools/cuda-ab/README.md
new file mode 100644
index 0000000..db87149
--- /dev/null
+++ b/tools/cuda-ab/README.md
@@ -0,0 +1,32 @@
+# CUDA A/B harness
+
+Run one deterministic CUDA workload on the **host** (bare-metal GPU, the
+oracle) and in the **guest** (Open-IOV emulated vGPU), then compare the
+machine-readable results. This mirrors the earlier host/guest `nvidia-smi`
+comparison, but exercises actual execution: allocation, HtoD/DtoH copies,
+kernel launch, synchronisation and verification.
+
+```
+tools/cuda-ab/ab.sh all       # build + host run + guest run + compare
+tools/cuda-ab/ab.sh host      # host only  -> run/host.txt
+tools/cuda-ab/ab.sh guest     # push binary, run in the running VM -> run/guest.txt
+tools/cuda-ab/ab.sh compare   # diff the stable result lines
+```
+
+## Files
+
+| File | Purpose |
+|---|---|
+| `vector_add.cu` | the workload (`DEVICE`/`RESULT`/`STATUS` output lines) |
+| `ab.sh` | host build/run, guest push/run, comparison |
+| `run/` | generated outputs (`host.txt`, `guest.txt`) |
+
+## Notes
+
+* The binary links the CUDA runtime **statically** (`--cudart=static`), so the
+  guest needs no CUDA toolkit - only the guest driver's `libcuda.so`.
+* The guest transfer uses the guest agent (base64 chunks); the VM must be
+  running (`scripts/vgpu_vm.py`).
+* `DEVICE` lines are normalised before diffing (the guest reports the
+  emulated GPU name); `RESULT` and `STATUS` must match exactly.
+* Host requirement: `nvidia-cuda-toolkit` (`nvcc`).
diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh
new file mode 100755
index 0000000..07f6b24
--- /dev/null
+++ b/tools/cuda-ab/ab.sh
@@ -0,0 +1,88 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# ab.sh - CUDA A/B harness: run the same workload on the host (bare metal,
+# real Turing GPU) and in the guest (Open-IOV emulated vGPU), then compare the
+# machine-readable results. Mirrors the nvidia-smi host/guest comparison used
+# earlier: the host run is the oracle.
+#
+# Usage:
+#   tools/cuda-ab/ab.sh host     # build + run on the host, save run/host.txt
+#   tools/cuda-ab/ab.sh guest    # push the binary + run in the running VM
+#   tools/cuda-ab/ab.sh compare  # diff the stable result lines
+#   tools/cuda-ab/ab.sh all      # host + guest + compare
+#
+# The binary links the CUDA runtime statically, so the guest needs no CUDA
+# toolkit - only the guest driver's libcuda.so.
+set -e
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+REPO="$(cd "$HERE/../.." && pwd)"
+VMEXEC="$REPO/scripts/vgpu_vm.py"
+RUN="$HERE/run"
+SRC="$HERE/vector_add.cu"
+BIN="$RUN/vector_add"
+
+mkdir -p "$RUN"
+
+build_host() {
+	command -v nvcc >/dev/null || { echo "nvcc not found (install nvidia-cuda-toolkit)"; exit 1; }
+	nvcc -O2 --cudart=static -o "$BIN" "$SRC"
+}
+
+run_host() {
+	build_host
+	echo "[*] host run (oracle)"
+	"$BIN" | tee "$RUN/host.txt"
+}
+
+push_bin() {
+	[ -x "$BIN" ] || build_host
+	local b64 chunk off total size
+	b64=$(base64 -w0 "$BIN")
+	total=${#b64}
+	off=0
+	size=70000
+	echo "[*] pushing $(stat -c%s "$BIN") bytes to guest"
+	"$VMEXEC" exec "rm -f /tmp/vadd.b64" >/dev/null 2>&1
+	while [ $off -lt $total ]; do
+		chunk=${b64:$off:$size}
+		"$VMEXEC" exec "echo -n '$chunk' >> /tmp/vadd.b64" >/dev/null 2>&1
+		off=$((off + size))
+	done
+	"$VMEXEC" exec "base64 -d /tmp/vadd.b64 > /tmp/vector_add && chmod +x /tmp/vector_add && rm -f /tmp/vadd.b64 && echo PUSH_OK"
+}
+
+run_guest() {
+	push_bin
+	echo "[*] guest run (emulated vGPU)"
+	"$VMEXEC" exec "/tmp/vector_add" 2>&1 | tee "$RUN/guest.txt"
+}
+
+compare() {
+	local h g
+	[ -f "$RUN/host.txt" ] || { echo "missing $RUN/host.txt (run host first)"; exit 1; }
+	[ -f "$RUN/guest.txt" ] || { echo "missing $RUN/guest.txt (run guest first)"; exit 1; }
+	h=$(grep -E '^(DEVICE|RESULT|STATUS)' "$RUN/host.txt")
+	g=$(grep -E '^(DEVICE|RESULT|STATUS)' "$RUN/guest.txt")
+	echo "--- host:";  echo "$h"
+	echo "--- guest:"; echo "$g"
+	# DEVICE lines carry the GPU name (differs by design); compare the rest.
+	h_norm=$(echo "$h" | sed 's/^DEVICE .*/DEVICE <gpu>/')
+	g_norm=$(echo "$g" | sed 's/^DEVICE .*/DEVICE <gpu>/')
+	if [ "$h_norm" = "$g_norm" ]; then
+		echo "A/B: MATCH"
+	else
+		echo "A/B: MISMATCH"
+		diff <(echo "$h_norm") <(echo "$g_norm") || true
+		exit 1
+	fi
+}
+
+case "${1:-all}" in
+host)    run_host ;;
+guest)   run_guest ;;
+compare) compare ;;
+all)     run_host; run_guest; compare ;;
+*)       echo "usage: $0 [host|guest|compare|all]"; exit 1 ;;
+esac
diff --git a/tools/cuda-ab/run/host.txt b/tools/cuda-ab/run/host.txt
new file mode 100644
index 0000000..53e7e38
--- /dev/null
+++ b/tools/cuda-ab/run/host.txt
@@ -0,0 +1,3 @@
+DEVICE NVIDIA GeForce RTX 2080 Super with Max-Q Design cc=7.5 sm=48 mem=8165523456
+RESULT n=1048576 bad=0 checksum=392344090.500000
+STATUS OK
diff --git a/tools/cuda-ab/run/vector_add b/tools/cuda-ab/run/vector_add
new file mode 100755
index 0000000..c1581c1
--- /dev/null
+++ b/tools/cuda-ab/run/vector_add
diff --git a/tools/cuda-ab/vector_add.cu b/tools/cuda-ab/vector_add.cu
new file mode 100644
index 0000000..2e3f1d6
--- /dev/null
+++ b/tools/cuda-ab/vector_add.cu
@@ -0,0 +1,93 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * vector_add.cu - CUDA A/B workload for the Open-IOV harness.
+ *
+ * One deterministic, machine-readable workload that can run on the host (bare
+ * metal) and in a guest (emulated vGPU). The output lines are stable so the
+ * two runs can be diffed:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   RESULT n=<count> bad=<mismatches> checksum=<sum>
+ *   STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation.
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cuda_runtime.h>
+
+#define CHECK(x)                                                         \
+	do {                                                             \
+		cudaError_t e_ = (x);                                    \
+		if (e_ != cudaSuccess) {                                 \
+			printf("CUDA_ERROR %s at line %d\n",             \
+			       cudaGetErrorString(e_), __LINE__);        \
+			printf("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+__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];
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int n = 1 << 20;
+	size_t bytes = (size_t)n * sizeof(float);
+	float *h_a, *h_b, *h_c;
+	float *d_a, *d_b, *d_c;
+	double sum = 0.0;
+	int bad = 0, i;
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount,
+	       (unsigned long long)prop.totalGlobalMem);
+
+	h_a = (float *)malloc(bytes);
+	h_b = (float *)malloc(bytes);
+	h_c = (float *)malloc(bytes);
+	if (!h_a || !h_b || !h_c)
+		return 1;
+
+	for (i = 0; i < n; i++) {
+		h_a[i] = (float)(i % 1000) * 0.5f;
+		h_b[i] = (float)(i % 997) * 0.25f;
+	}
+
+	CHECK(cudaMalloc(&d_a, bytes));
+	CHECK(cudaMalloc(&d_b, bytes));
+	CHECK(cudaMalloc(&d_c, bytes));
+	CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
+	CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
+
+	vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
+
+	for (i = 0; i < n; i++) {
+		if (h_c[i] != h_a[i] + h_b[i])
+			bad++;
+		sum += h_c[i];
+	}
+
+	printf("RESULT n=%d bad=%d checksum=%.6f\n", n, bad, sum);
+	printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
+
+	CHECK(cudaFree(d_a));
+	CHECK(cudaFree(d_b));
+	CHECK(cudaFree(d_c));
+	free(h_a);
+	free(h_b);
+	free(h_c);
+	return bad == 0 ? 0 : 1;
+}