// SPDX-License-Identifier: BSD-3-Clause
/*
 * r18_anim.cu - GPU-computed animation for R18 (v3: single-page kernel).
 *
 * Renders NFRAMES frames of an animated plasma with one launch per frame.
 * The emulated forward path delivers the kernel through ONE 4 KB code page
 * (the window shadow copies the program page; multi-page kernels execute off
 * the end of the mapped page - the 904-instruction v2 kernel faulted with
 * "Illegal Instruction Encoding" on the real engine). This version keeps the
 * SASS text under 4096 bytes: 512x360 RGBA8 (power-of-two width for cheap
 * index math), grid-stride loop, three MUFU.SIN under --use_fast_math.
 *
 * ONE 0.9 MB device frame buffer is reused for every frame (small-alloc
 * regime, matching the proven CUDA path); after each launch the frame is
 * copied back (DtoH) and APPENDED to /dev/shm with a magic header so the
 * host can find it by scanning the guest RAM file. The file is
 *
 *   "R18FRM01" u32 count u32 width u32 height u32 stride
 *   then count*height*stride bytes of RGBA8
 *
 * The header is written first with the final count and bytes are appended
 * (flushed) per frame, so a host-side harvest can recover every completed
 * frame even if the client or guest wedges mid-run.
 *
 * Build (host):
 *   nvcc -O2 --use_fast_math --cudart=static -arch=sm_75 -o r18_anim r18_anim.cu
 * Usage: r18_anim [frames] [outpath] [blocks]
 */

#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cstring>
#include <ctime>
#include <cuda_runtime.h>

#define WIDTH   512
#define HEIGHT  360
#define NFRAMES 60
#define OUTPATH "/dev/shm/r18frames.bin"

#define CHECK(x)                                                         \
	do {                                                             \
		cudaError_t e_ = (x);                                    \
		if (e_ != cudaSuccess) {                                 \
			printf("CUDA_ERROR %s at line %d\n",             \
			       cudaGetErrorString(e_), __LINE__);        \
			printf("STATUS FAIL\n");                         \
			return 1;                                        \
		}                                                        \
	} while (0)

/* 3 MUFU.SIN only - the text must stay under one 4 KB page. */
__global__ void plasma_kernel(uchar4 *fb, int w, int h, int frame)
{
	unsigned stride = gridDim.x * blockDim.x;
	unsigned xmask = (unsigned)w - 1u;	/* w is a power of two */
	int idx;

	for (idx = blockIdx.x * blockDim.x + threadIdx.x; idx < w * h;
	     idx += stride) {
		float u = (float)((unsigned)idx & xmask) * (1.0f / 512.0f);
		float v = (float)((unsigned)idx >> 9) * (1.0f / 360.0f);
		float t = (float)frame * 0.10f;
		float p = sinf(u * 9.0f + t * 2.0f)
			+ sinf(v * 7.0f - t * 2.0f);
		float s = sinf(p * 1.5f);

		p *= 0.5f;		/* -1 .. 1 */
		fb[idx] = make_uchar4((unsigned char)(127.5f + 127.5f * s),
				      (unsigned char)(127.5f + 127.5f * p),
				      (unsigned char)(127.5f - 127.5f * s),
				      255);
	}
}

static uint64_t fnv1a(const unsigned char *p, size_t n)
{
	uint64_t h = 1469598103934665603ull;

	while (n--) {
		h ^= *p++;
		h *= 1099511628211ull;
	}
	return h;
}

static double now_ms(void)
{
	struct timespec ts;

	clock_gettime(CLOCK_MONOTONIC, &ts);
	return (double)ts.tv_sec * 1e3 + (double)ts.tv_nsec / 1e6;
}

int main(int argc, char **argv)
{
	int nframes = NFRAMES;
	const char *outpath = OUTPATH;
	size_t fbytes = (size_t)WIDTH * HEIGHT * 4;
	uchar4 *d_fb = NULL;
	unsigned char *h_fb = NULL;
	double km = 0.0, cm = 0.0, t0, t1, wall0 = now_ms();
	int blocks = 720;
	int f, dev = -1;
	FILE *fp;
	uint32_t hdr[4];

	if (argc > 1)
		nframes = atoi(argv[1]);
	if (nframes <= 0 || nframes > 60000)
		nframes = NFRAMES;
	if (argc > 2)
		outpath = argv[2];
	if (argc > 3)
		blocks = atoi(argv[3]);
	if (blocks <= 0 || blocks > 4096)
		blocks = 720;

	setvbuf(stdout, NULL, _IONBF, 0);

	CHECK(cudaGetDevice(&dev));
	{
		cudaDeviceProp prop;

		if (cudaGetDeviceProperties(&prop, dev) == cudaSuccess)
			printf("DEVICE %s cc=%d.%d\n", prop.name,
			       prop.major, prop.minor);
	}
	printf("ANIM start frames=%d %dx%d blocks=%d out=%s\n",
	       nframes, WIDTH, HEIGHT, blocks, outpath);

	CHECK(cudaMalloc(&d_fb, fbytes));
	h_fb = (unsigned char *)malloc(fbytes);
	if (!h_fb) {
		printf("HOST_ALLOC_FAIL\nSTATUS FAIL\n");
		return 1;
	}

	fp = fopen(outpath, "wb");
	if (!fp) {
		printf("OPEN_FAIL %s\nSTATUS FAIL\n", outpath);
		return 1;
	}
	hdr[0] = (uint32_t)nframes;
	hdr[1] = WIDTH;
	hdr[2] = HEIGHT;
	hdr[3] = WIDTH * 4;
	if (fwrite("R18FRM01", 1, 8, fp) != 8 ||
	    fwrite(hdr, 1, sizeof(hdr), fp) != sizeof(hdr) || fflush(fp)) {
		printf("WRITE_FAIL header\nSTATUS FAIL\n");
		fclose(fp);
		return 1;
	}

	for (f = 0; f < nframes; f++) {
		uchar4 *dst = d_fb;
		uint64_t sum = 0;
		size_t i;

		t0 = now_ms();
		plasma_kernel<<<blocks, 256>>>(dst, WIDTH, HEIGHT, f);
		CHECK(cudaGetLastError());
		CHECK(cudaDeviceSynchronize());
		t1 = now_ms();
		km += t1 - t0;
		CHECK(cudaMemcpy(h_fb, dst, fbytes, cudaMemcpyDeviceToHost));
		cm += now_ms() - t1;

		for (i = 0; i < fbytes; i++)
			sum += h_fb[i];
		if (fwrite(h_fb, 1, fbytes, fp) != fbytes || fflush(fp)) {
			printf("WRITE_FAIL frame %d\nSTATUS FAIL\n", f);
			fclose(fp);
			return 1;
		}
		printf("FRAME %03d fnv=0x%016llx mean=%u kms=%.1f cms=%.1f wall=%.1f\n",
		       f, (unsigned long long)fnv1a(h_fb, fbytes),
		       (unsigned)(sum / fbytes), km, cm, now_ms() - wall0);
	}

	fclose(fp);
	printf("KERNEL frames=%d %dx%d blocks=%d launches=%d total_ms=%.2f per_frame_ms=%.2f\n",
	       nframes, WIDTH, HEIGHT, blocks, nframes, km, km / nframes);
	printf("COPY frames=%d total_ms=%.2f per_frame_ms=%.2f\n", nframes,
	       cm, cm / nframes);
	printf("WRITE path=%s bytes=%zu magic=R18FRM01 count=%d %dx%d stride=%d wall_ms=%.1f\n",
	       outpath, 24 + fbytes * (size_t)nframes, nframes, WIDTH, HEIGHT,
	       WIDTH * 4, now_ms() - wall0);
	printf("STATUS OK\n");

	CHECK(cudaFree(d_fb));
	free(h_fb);
	return 0;
}