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, 40 insertions, 0 deletions

diff --git a/tools/cuda-bridge/kernels.cu b/tools/cuda-bridge/kernels.cu
new file mode 100644
index 0000000..508c94c
--- /dev/null
+++ b/tools/cuda-bridge/kernels.cu
@@ -0,0 +1,40 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// kernels.cu - CUDA kernels for the host-side CUDA bridge.
+//
+// Exactly mirrors the victory_mail semantics: mail_kernel writes every
+// character of the message from device memory, vadd computes c = a + b.
+// vadd_mutated is the falsifiability variant (adds 1.0f); it exists only to
+// prove the bridge executes the real requested kernel, never a canned answer.
+//
+// Compiled to PTX at build time by nvcc (see Makefile); the PTX is embedded
+// into cuda_bridge and loaded with cuModuleLoadData.
+
+#define MAIL "the little gpu man says you've got mail"
+
+__device__ const char kMail[] = MAIL;
+
+extern "C" __global__ void mail_kernel(char *out, int len)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < len)
+		out[i] = kMail[i];
+}
+
+extern "C" __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];
+}
+
+extern "C" __global__ void vadd_mutated(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] + 1.0f;
+}