cuda-ab: launch_count - single-launch execution-count repro (each forwarded launch runs its engine work twice; the multi-pass class root cause)

1 files changed, 25 insertions, 0 deletions

diff --git a/tools/cuda-ab/launch_count.cu b/tools/cuda-ab/launch_count.cu
new file mode 100644
index 0000000..02db1d6
--- /dev/null
+++ b/tools/cuda-ab/launch_count.cu
@@ -0,0 +1,25 @@
+// SPDX-License-Identifier: BSD-3-Clause
+// launch_count.cu - single-launch execution-count repro (r342).
+//
+// One launch, one atomic increment. The engine-written page witness showed the
+// increment applied TWICE (1 -> 2; the 100-increment variant -> 200): every
+// forwarded launch executes its engine work twice (one submission, one ring
+// publish, completion OK). This is the root cause of the multi-pass content
+// class (victory_reduce exactly 2x, scale/saxpy/smiley/multikernel wrong).
+// Run in the guest; expected LAUNCH_COUNT=1.
+// for a single launch, the launch path executes work more than once.
+#include <cstdio>
+#include <cuda_runtime.h>
+__global__ void inc(unsigned long long *c) { if (threadIdx.x == 0 && blockIdx.x == 0) atomicAdd(c, 1ULL); }
+int main(void) {
+	unsigned long long *d; unsigned long long h = 0;
+	if (cudaMalloc(&d, 8) != cudaSuccess) { printf("CUDA_ERROR malloc\nSTATUS FAIL\n"); return 1; }
+	if (cudaMemset(d, 0, 8) != cudaSuccess) { printf("CUDA_ERROR memset\nSTATUS FAIL\n"); return 1; }
+	inc<<<1, 32>>>(d);
+	if (cudaGetLastError() != cudaSuccess) { printf("CUDA_ERROR launch\nSTATUS FAIL\n"); return 1; }
+	if (cudaDeviceSynchronize() != cudaSuccess) { printf("CUDA_ERROR sync\nSTATUS FAIL\n"); return 1; }
+	if (cudaMemcpy(&h, d, 8, cudaMemcpyDeviceToHost) != cudaSuccess) { printf("CUDA_ERROR d2h\nSTATUS FAIL\n"); return 1; }
+	printf("LAUNCH_COUNT=%llu (expect 1)\n", h);
+	printf("STATUS %s\n", h == 1 ? "OK" : "FAIL");
+	return h == 1 ? 0 : 1;
+}