// SPDX-License-Identifier: BSD-3-Clause
/*
 * p2_managed_probe - EXPECTED-FAILURE PROBE (documentation, not acceptance).
 *
 * Managed (UVM) memory is a known unsupported path for the emulated vGPU.
 * This probe attempts cudaMallocManaged, records the outcome as a PROBE
 * line, clears any error, and then proves that ordinary cudaMalloc + launch
 * still works afterwards. The emulated guest is expected to report a
 * failure here; the host oracle may report success. p2/run_guest.sh treats
 * this program as informational (content not diffed), but requires the
 * program to complete with STATUS OK - i.e. cleanly report the managed
 * behaviour and keep working - so a hang, crash or poisoned context still
 * fails the suite.
 *
 * Device memory: 4 KB managed (if it works) + 4 KB normal.
 */
#include "common.h"

__global__ void fillk(int *p, int v, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < n)
		p[i] = v + i;
}

int main(int argc, char **argv)
{
	if (p2_emit_mode(argc, argv)) {
		printf("PROBE managed state=no-oracle-gpu note=emit-mode\n");
		printf("STATUS OK\n");
		return 0;
	}

	p2_device_info();

	int *p = NULL;
	cudaError_t e = cudaMallocManaged(&p, sizeof(int) * 4);

	if (e == cudaSuccess) {
		cudaError_t e2, e3;
		int ok = 0;

		fillk<<<1, 1>>>(p, 5, 4);
		e2 = cudaGetLastError();
		e3 = cudaDeviceSynchronize();
		if (e2 == cudaSuccess && e3 == cudaSuccess)
			ok = (p[0] == 5 && p[1] == 6 && p[2] == 7 &&
			      p[3] == 8) ? 1 : 0;
		printf("PROBE managed alloc=ok write_read=%s\n",
		       ok ? "ok" : "bad");
		cudaFree(p);
	} else {
		printf("PROBE managed alloc=fail rc=%d name=%s\n", (int)e,
		       cudaGetErrorString(e));
		(void)cudaGetLastError();
	}

	int *d = NULL, bad = 0;
	cudaError_t e4 = cudaMalloc(&d, sizeof(int) * 4);

	if (e4 == cudaSuccess) {
		int h[4];

		fillk<<<1, 1>>>(d, 100, 4);
		(void)cudaGetLastError();
		cudaDeviceSynchronize();
		cudaMemcpy(h, d, sizeof(h), cudaMemcpyDeviceToHost);
		for (int i = 0; i < 4; i++)
			if (h[i] != 100 + i)
				bad++;
		cudaFree(d);
		printf("PROBE fallback rc=0 bad=%d\n", bad);
	} else {
		printf("PROBE fallback rc=%d name=%s\n", (int)e4,
		       cudaGetErrorString(e4));
		bad = 1;
	}

	return p2_status(bad);
}