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;
+}