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

diff --git a/tools/cuda-bridge/kernels.ptx b/tools/cuda-bridge/kernels.ptx
new file mode 100644
index 0000000..4d1b3ad
--- /dev/null
+++ b/tools/cuda-bridge/kernels.ptx
@@ -0,0 +1,132 @@
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-32267302
+// Cuda compilation tools, release 12.0, V12.0.140
+// Based on NVVM 7.0.1
+//
+
+.version 8.0
+.target sm_75
+.address_size 64
+
+	// .globl	mail_kernel
+.global .align 1 .b8 kMail[40] = {116, 104, 101, 32, 108, 105, 116, 116, 108, 101, 32, 103, 112, 117, 32, 109, 97, 110, 32, 115, 97, 121, 115, 32, 121, 111, 117, 39, 118, 101, 32, 103, 111, 116, 32, 109, 97, 105, 108, 0};
+
+.visible .entry mail_kernel(
+	.param .u64 mail_kernel_param_0,
+	.param .u32 mail_kernel_param_1
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .b16 	%rs<3>;
+	.reg .b32 	%r<6>;
+	.reg .b64 	%rd<7>;
+
+
+	ld.param.u64 	%rd1, [mail_kernel_param_0];
+	ld.param.u32 	%r2, [mail_kernel_param_1];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ntid.x;
+	mov.u32 	%r5, %tid.x;
+	mad.lo.s32 	%r1, %r3, %r4, %r5;
+	setp.ge.s32 	%p1, %r1, %r2;
+	@%p1 bra 	$L__BB0_2;
+
+	cvta.to.global.u64 	%rd2, %rd1;
+	cvt.s64.s32 	%rd3, %r1;
+	mov.u64 	%rd4, kMail;
+	add.s64 	%rd5, %rd4, %rd3;
+	ld.global.nc.u8 	%rs1, [%rd5];
+	cvt.u16.u8 	%rs2, %rs1;
+	add.s64 	%rd6, %rd2, %rd3;
+	st.global.u8 	[%rd6], %rs2;
+
+$L__BB0_2:
+	ret;
+
+}
+	// .globl	vadd
+.visible .entry vadd(
+	.param .u64 vadd_param_0,
+	.param .u64 vadd_param_1,
+	.param .u64 vadd_param_2,
+	.param .u32 vadd_param_3
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .f32 	%f<4>;
+	.reg .b32 	%r<6>;
+	.reg .b64 	%rd<11>;
+
+
+	ld.param.u64 	%rd1, [vadd_param_0];
+	ld.param.u64 	%rd2, [vadd_param_1];
+	ld.param.u64 	%rd3, [vadd_param_2];
+	ld.param.u32 	%r2, [vadd_param_3];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ntid.x;
+	mov.u32 	%r5, %tid.x;
+	mad.lo.s32 	%r1, %r3, %r4, %r5;
+	setp.ge.s32 	%p1, %r1, %r2;
+	@%p1 bra 	$L__BB1_2;
+
+	cvta.to.global.u64 	%rd4, %rd1;
+	mul.wide.s32 	%rd5, %r1, 4;
+	add.s64 	%rd6, %rd4, %rd5;
+	cvta.to.global.u64 	%rd7, %rd2;
+	add.s64 	%rd8, %rd7, %rd5;
+	ld.global.f32 	%f1, [%rd8];
+	ld.global.f32 	%f2, [%rd6];
+	add.f32 	%f3, %f2, %f1;
+	cvta.to.global.u64 	%rd9, %rd3;
+	add.s64 	%rd10, %rd9, %rd5;
+	st.global.f32 	[%rd10], %f3;
+
+$L__BB1_2:
+	ret;
+
+}
+	// .globl	vadd_mutated
+.visible .entry vadd_mutated(
+	.param .u64 vadd_mutated_param_0,
+	.param .u64 vadd_mutated_param_1,
+	.param .u64 vadd_mutated_param_2,
+	.param .u32 vadd_mutated_param_3
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .f32 	%f<5>;
+	.reg .b32 	%r<6>;
+	.reg .b64 	%rd<11>;
+
+
+	ld.param.u64 	%rd1, [vadd_mutated_param_0];
+	ld.param.u64 	%rd2, [vadd_mutated_param_1];
+	ld.param.u64 	%rd3, [vadd_mutated_param_2];
+	ld.param.u32 	%r2, [vadd_mutated_param_3];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ntid.x;
+	mov.u32 	%r5, %tid.x;
+	mad.lo.s32 	%r1, %r3, %r4, %r5;
+	setp.ge.s32 	%p1, %r1, %r2;
+	@%p1 bra 	$L__BB2_2;
+
+	cvta.to.global.u64 	%rd4, %rd1;
+	mul.wide.s32 	%rd5, %r1, 4;
+	add.s64 	%rd6, %rd4, %rd5;
+	cvta.to.global.u64 	%rd7, %rd2;
+	add.s64 	%rd8, %rd7, %rd5;
+	ld.global.f32 	%f1, [%rd8];
+	ld.global.f32 	%f2, [%rd6];
+	add.f32 	%f3, %f2, %f1;
+	add.f32 	%f4, %f3, 0f3F800000;
+	cvta.to.global.u64 	%rd9, %rd3;
+	add.s64 	%rd10, %rd9, %rd5;
+	st.global.f32 	[%rd10], %f4;
+
+$L__BB2_2:
+	ret;
+
+}
+