vgsp: host CUDA bridge + sema release recovery; greeneye simulator wave

- kernel: gsp_cuda_bridge (QMD launch capture, debugfs job/out/done, param+output resolution cached from inline writes), sema drain fixes (never hold/starve a channel), GR ring autopin + counter autoreg, per-page push reads, counter-derive gate; module builds clean
- tools/cuda-bridge: CUDA driver-API daemon executing mail/vadd kernels on the physical GPU, CUBB2/CBD2 v2 protocol, layout check + self-tests
- tools/guest-diag + tools/cuda-ab: scanner/pin fixes, victory push-only mode
- debug/risc-v-gsp-rm-vm: Phase S simulator wave (GSP device model, boot ladder, triage/re-demo harnesses, docs 30-41), GPGPU-sim scaffolding
- STATE-CURRENT/docs updates

1 files changed, 152 insertions, 0 deletions

diff --git a/tools/cuda-ab/victory_mail.cu b/tools/cuda-ab/victory_mail.cu
new file mode 100644
index 0000000..e83fc1b
--- /dev/null
+++ b/tools/cuda-ab/victory_mail.cu
@@ -0,0 +1,152 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * victory_mail.cu - the Open-IOV victory program.
+ *
+ * Run from the guest. Two kernels execute on the physical GPU:
+ *
+ *   mail_kernel  - the GPU itself writes every character of the message into
+ *                  device memory; the host only copies the bytes back.
+ *   vadd         - the classic vectorAdd acceptance check, verified element
+ *                  by element against the host-computed result.
+ *
+ * The host prints the message only after it has been authored by GPU
+ * execution and copied back byte-for-byte, so the line below is proof that a
+ * kernel really ran on the device:
+ *
+ *   the little gpu man says you've got mail
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   MAIL <message>
+ *   VADD 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:
+ *
+ *   nvcc -O2 --cudart=static -o victory_mail victory_mail.cu
+ *
+ * Then, from the guest:  /tmp/victory_mail
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cuda_runtime.h>
+
+#define MAIL "the little gpu man says you've got mail"
+
+#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)
+
+/* The message lives in device code: the GPU both provides and writes it. */
+__device__ const char kMail[] = MAIL;
+
+/* Each thread writes exactly one character of the message. */
+__global__ void mail_kernel(char *out, int len)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < len)
+		out[i] = kMail[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];
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int len = (int)strlen(MAIL);
+	const int n = 1 << 20;
+	size_t bytes = (size_t)n * sizeof(float);
+	char *d_mail, h_mail[128];
+	float *h_a, *h_b, *h_c;
+	float *d_a, *d_b, *d_c;
+	double sum = 0.0;
+	int bad = 0, i;
+
+	/* Unbuffered: progress must be visible even if a later call blocks. */
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	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);
+
+	/* The GPU writes the mail. */
+	CHECK(cudaMalloc(&d_mail, len));
+	mail_kernel<<<1, 64>>>(d_mail, len);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	memset(h_mail, 0, sizeof(h_mail));
+	CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost));
+	if (memcmp(h_mail, MAIL, len) != 0) {
+		printf("MAIL bad=<kernel output did not match>\n");
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+	printf("MAIL %s\n", h_mail);
+
+	/* The classic acceptance check: vectorAdd. */
+	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("VADD n=%d bad=%d checksum=%.6f\n", n, bad, sum);
+
+	if (bad != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	printf("the little gpu man says you've got mail\n");
+	printf("STATUS OK\n");
+
+	CHECK(cudaFree(d_mail));
+	CHECK(cudaFree(d_a));
+	CHECK(cudaFree(d_b));
+	CHECK(cudaFree(d_c));
+	free(h_a);
+	free(h_b);
+	free(h_c);
+	return 0;
+}