// SPDX-License-Identifier: BSD-3-Clause
/*
 * p2_printf_err - device printf, a passing device assert, and an intentional
 * launch error path.
 *
 *   1. A 4-thread kernel prints deterministic P2PRINT lines from the device.
 *   2. A kernel with assert(n == 4) exercises the device-side assert
 *      machinery on its passing path (a fired assert is deliberately not
 *      tested: it traps the kernel and poisons the context).
 *   3. A kernel launched with gridDim.x == 0 must be rejected:
 *      cudaErrorInvalidConfiguration (9), and the error must be cleared by
 *      cudaGetLastError. A blockDim.x == 0 launch is exercised too.
 *   4. A normal launch after the errors proves error recovery.
 *
 * Expected error strings come from the statically linked CUDA runtime, so
 * they must be identical on host and guest.
 * Device memory: out 4 B.
 */
#include "common.h"
#include <cassert>

__global__ void hello(int nthreads)
{
	int i = threadIdx.x;

	if (i < nthreads)
		printf("P2PRINT hello i=%d\n", i);
}

__global__ void second(void)
{
	if (threadIdx.x == 0)
		printf("P2PRINT second c=%c v=%d\n", 'Z', 7);
}

__global__ void simple(int *out)
{
	out[0] = 42;
}

__global__ void assertk(int *out, int n)
{
	assert(n == 4);
	if (threadIdx.x == 0)
		out[0] = n;
}

int main(int argc, char **argv)
{
	if (p2_emit_mode(argc, argv)) {
		printf("P2PRINT hello i=0\n");
		printf("P2PRINT hello i=1\n");
		printf("P2PRINT hello i=2\n");
		printf("P2PRINT hello i=3\n");
		printf("P2PRINT second c=Z v=7\n");
		printf("ERR launch grid0 rc=%d name=%s\n",
		       (int)cudaErrorInvalidConfiguration,
		       cudaGetErrorString(cudaErrorInvalidConfiguration));
		printf("ERR launch block0 rejected=1\n");
		printf("ERR cleared rc=0\n");
		printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n");
		printf("RESULT assert n=4 bad=0\n");
		return p2_status(0);
	}

	p2_device_info();

	int *d_out = NULL;
	P2_CHECK(cudaMalloc(&d_out, sizeof(int)));

	hello<<<1, 4>>>(4);
	P2_CHECK(cudaGetLastError());
	P2_CHECK(cudaDeviceSynchronize());
	second<<<1, 4>>>();
	P2_CHECK(cudaGetLastError());
	P2_CHECK(cudaDeviceSynchronize());

	(void)cudaGetLastError();

	int badgrid = 0, badblock = 0;
	cudaError_t e_grid, e_block, cleared;
	int got = 0;

	simple<<<badgrid, 1>>>(d_out);
	e_grid = cudaGetLastError();
	printf("ERR launch grid0 rc=%d name=%s\n", (int)e_grid,
	       cudaGetErrorString(e_grid));

	simple<<<1, badblock>>>(d_out);
	e_block = cudaGetLastError();
	printf("ERR launch block0 rejected=%d\n",
	       (e_block != cudaSuccess) ? 1 : 0);

	cleared = cudaGetLastError();
	printf("ERR cleared rc=%d\n", (int)cleared);

	simple<<<1, 1>>>(d_out);
	P2_CHECK(cudaGetLastError());
	P2_CHECK(cudaDeviceSynchronize());
	P2_CHECK(cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost));

	int grid0_bad = (e_grid != cudaErrorInvalidConfiguration) ? 1 : 0;
	int block0_bad = (e_block == cudaSuccess) ? 1 : 0;
	int clear_bad = (cleared != cudaSuccess) ? 1 : 0;
	int recovered = (got == 42) ? 1 : 0;
	int assert_bad = 0;

	assertk<<<1, 1>>>(d_out, 4);
	P2_CHECK(cudaGetLastError());
	P2_CHECK(cudaDeviceSynchronize());
	{
		int agot = 0;

		P2_CHECK(cudaMemcpy(&agot, d_out, sizeof(int),
				    cudaMemcpyDeviceToHost));
		assert_bad = (agot == 4) ? 0 : 1;
	}

	int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1) +
		  assert_bad;

	printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n",
	       grid0_bad, block0_bad, clear_bad, recovered);
	printf("RESULT assert n=4 bad=%d\n", assert_bad);

	cudaFree(d_out);
	return p2_status(bad);
}