// SPDX-License-Identifier: BSD-3-Clause
/*
 * p2_churn - cudaMalloc/cudaFree churn between launches.
 *
 * Six iterations allocate 1,2,3,4,1,2 buffers of 2048 ints (8 KB each),
 * fill each with a seed-derived kernel, verify, then free all of them.
 * A final fresh allocation + launch proves the allocator remains healthy.
 * Peak live device memory: 32 KB (4 pages); 14 allocations overall.
 */
#include "common.h"

#define CH_N 2048
#define ITERS 6
#define FINAL 7

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

	if (i < n)
		p[i] = seed * 1000 + i;
}

int main(int argc, char **argv)
{
	int *got = (int *)malloc(sizeof(int) * CH_N);
	int bad = 0;

	if (!got)
		return 1;

	if (p2_emit_mode(argc, argv)) {
		for (int it = 0; it < ITERS; it++) {
			int nbuf = 1 + (it % 4);

			printf("RESULT churn iter=%d nbuf=%d n=%d bad=0\n", it,
			       nbuf, CH_N);
		}
		printf("RESULT churn final_buf=1 n=%d bad=0\n", CH_N);
		printf("RESULT churn iters=%d allocs=%d bad=0\n", ITERS,
		       1 + 2 + 3 + 4 + 1 + 2 + 1);
		free(got);
		return p2_status(0);
	}

	p2_device_info();

	for (int it = 0; it < ITERS; it++) {
		int nbuf = 1 + (it % 4);
		int *d[4] = { NULL, NULL, NULL, NULL };
		int sbad = 0, seed = it + 1;

		for (int k = 0; k < nbuf; k++) {
			P2_CHECK(cudaMalloc(&d[k], sizeof(int) * CH_N));
			fill<<<(CH_N + 255) / 256, 256>>>(d[k], seed, CH_N);
			P2_CHECK(cudaGetLastError());
			P2_CHECK(cudaDeviceSynchronize());
			P2_CHECK(cudaMemcpy(got, d[k], sizeof(int) * CH_N,
					    cudaMemcpyDeviceToHost));
			for (int i = 0; i < CH_N; i++)
				if (got[i] != seed * 1000 + i)
					sbad++;
		}
		for (int k = 0; k < nbuf; k++)
			P2_CHECK(cudaFree(d[k]));
		printf("RESULT churn iter=%d nbuf=%d n=%d bad=%d\n", it, nbuf,
		       CH_N, sbad);
		bad += sbad;
	}

	{
		int *d = NULL;
		int fbad = 0;

		P2_CHECK(cudaMalloc(&d, sizeof(int) * CH_N));
		fill<<<(CH_N + 255) / 256, 256>>>(d, FINAL, CH_N);
		P2_CHECK(cudaGetLastError());
		P2_CHECK(cudaDeviceSynchronize());
		P2_CHECK(cudaMemcpy(got, d, sizeof(int) * CH_N,
				    cudaMemcpyDeviceToHost));
		for (int i = 0; i < CH_N; i++)
			if (got[i] != FINAL * 1000 + i)
				fbad++;
		P2_CHECK(cudaFree(d));
		printf("RESULT churn final_buf=1 n=%d bad=%d\n", CH_N, fbad);
		bad += fbad;
	}

	cudaError_t fin = cudaGetLastError();

	if (fin != cudaSuccess) {
		printf("ERR churn final_rc=%d name=%s\n", (int)fin,
		       cudaGetErrorString(fin));
		bad++;
	}
	printf("RESULT churn iters=%d allocs=%d bad=%d\n", ITERS,
	       1 + 2 + 3 + 4 + 1 + 2 + 1, bad);
	free(got);
	return p2_status(bad);
}