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