// SPDX-License-Identifier: BSD-3-Clause
/*
 * smiley.cu - vectorAdd draws a smiley face.
 *
 * Four vectorAdd passes execute on the GPU; each computes one scalar field
 *
 *   c[i] = a[i] + b[i]
 *
 * over the same n = 32768 elements laid out as a 256x128 sample grid:
 *
 *   pass 0  a=u^2          b=v^2          -> face disk field
 *   pass 1  a=(u+xe)^2     b=(v+ye)^2     -> left-eye disk field
 *   pass 2  a=(u-xe)^2     b=(v+ye)^2     -> right-eye disk field
 *   pass 3  a=u^2          b=(v-ym)^2     -> mouth annulus field
 *
 * with u,v centre-relative, aspect-corrected coordinates. The host copies
 * the four fields back and thresholds them to render a 64x32 ASCII smiley
 * (4x4 samples per character cell). The artwork therefore comes out of GPU
 * addition results only - no GPU-side drawing, no CPU-side math on pixels
 * beyond the thresholds.
 *
 * Output lines are stable and greppable (matches tools/cuda-ab/ab.sh):
 *
 *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
 *   MALLOC bytes=<n> rc=<0|err>
 *   VADD passes=4 n=<count> bad=<mismatches> checksum=<sum>
 *   RESULT cells=<W>x<H> inked=<cells> checksum=<sum>
 *   STATUS OK|FAIL
 *
 * Build with a static CUDA runtime (--cudart=static) so the guest only needs
 * the driver's libcuda.so, not a CUDA toolkit installation:
 *
 *   nvcc -O2 --cudart=static -o smiley smiley.cu
 *
 * Then, from the guest:  /tmp/smiley
 */

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

#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)

__global__ void vadd(const float *a, const float *b, float *c, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < n)
		c[i] = a[i] + b[i];
}

/* sample grid */
#define GW 256
#define GH 128
#define GN (GW * GH)

/* character cell = 4x4 samples */
#define CW 64
#define CH 32
#define CELL 4

/* feature geometry, in aspect-corrected grid units (u ~ +/-127, v ~ +/-127) */
#define FACE_R2   10000.0f /* face disk:    u^2+v^2   <= 100^2   */
#define EYE_X     36.0f
#define EYE_Y     30.0f
#define EYE_R2    144.0f   /* eyes: r=12 */
#define MOUTH_Y   24.0f
#define MOUTH_IN  2025.0f  /* 45^2 */
#define MOUTH_OUT 3025.0f  /* 55^2 */
#define MOUTH_V   40.0f

int main(void)
{
	cudaDeviceProp prop;
	float *h_a, *h_b, *h_out, *fields;
	float *d_a, *d_b, *d_c;
	size_t bytes = (size_t)GN * sizeof(float);
	double checksum = 0.0;
	int bad = 0, pass, i;

	setvbuf(stdout, NULL, _IONBF, 0);

	CHECK(cudaGetDeviceProperties(&prop, 0));
	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
	       prop.major, prop.minor, prop.multiProcessorCount,
	       (unsigned long long)prop.totalGlobalMem);

	{
		cudaError_t e_a = cudaMalloc(&d_a, bytes);
		cudaError_t e_b = e_a == cudaSuccess ?
			cudaMalloc(&d_b, bytes) : e_a;
		cudaError_t e_c = e_b == cudaSuccess ?
			cudaMalloc(&d_c, bytes) : e_b;

		printf("MALLOC bytes=%zu rc=%d\n", bytes, (int)e_a);
		if (e_a != cudaSuccess) {
			size_t fm = 0, tm = 0;

			cudaMemGetInfo(&fm, &tm);
			printf("MEMINFO free=%llu total=%llu\n",
			       (unsigned long long)fm,
			       (unsigned long long)tm);
		}
		CHECK(e_a);
		CHECK(e_b);
		CHECK(e_c);
	}

	h_a = (float *)malloc(bytes);
	h_b = (float *)malloc(bytes);
	h_out = (float *)malloc(bytes);
	fields = (float *)malloc((size_t)4 * bytes);
	if (!h_a || !h_b || !h_out || !fields)
		return 1;

	for (pass = 0; pass < 4; pass++) {
		for (i = 0; i < GN; i++) {
			int x = i % GW, y = i / GW;
			float u = (float)x - 127.5f;
			float v = ((float)y - 63.5f) * 2.0f;
			float pa, pb;

			switch (pass) {
			case 0: /* face */
				pa = u * u;            pb = v * v;
				break;
			case 1: /* left eye */
				pa = (u + EYE_X) * (u + EYE_X);
				pb = (v + EYE_Y) * (v + EYE_Y);
				break;
			case 2: /* right eye */
				pa = (u - EYE_X) * (u - EYE_X);
				pb = (v + EYE_Y) * (v + EYE_Y);
				break;
			default: /* mouth */
				pa = u * u;
				pb = (v - MOUTH_Y) * (v - MOUTH_Y);
				break;
			}
			h_a[i] = pa;
			h_b[i] = pb;
		}

		CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
		CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
		vadd<<<(GN + 255) / 256, 256>>>(d_a, d_b, d_c, GN);
		CHECK(cudaGetLastError());
		CHECK(cudaDeviceSynchronize());
		CHECK(cudaMemcpy(h_out, d_c, bytes, cudaMemcpyDeviceToHost));

		for (i = 0; i < GN; i++) {
			if (h_out[i] != h_a[i] + h_b[i])
				bad++;
			checksum += (double)h_out[i];
		}
		for (i = 0; i < GN; i++)
			fields[(size_t)pass * GN + i] = h_out[i];
	}

	printf("VADD passes=4 n=%d bad=%d checksum=%.0f\n", GN, bad, checksum);
	if (bad != 0) {
		printf("STATUS FAIL\n");
		return 1;
	}

	{
		static char art[CH][CW + 1];
		int inked = 0, cy, cx;

		for (cy = 0; cy < CH; cy++) {
			for (cx = 0; cx < CW; cx++) {
				int hit = 0, sy, sx;
				const float *f0 = fields;
				const float *f1 = fields + GN;
				const float *f2 = fields + 2 * GN;
				const float *f3 = fields + 3 * GN;

				for (sy = 0; sy < CELL; sy++) {
					int y = cy * CELL + sy;
					float v = ((float)y - 63.5f) * 2.0f;

					for (sx = 0; sx < CELL; sx++) {
						int x = cx * CELL + sx;
						int idx = y * GW + x;
						int inside;

						inside = f0[idx] <= FACE_R2;
						if (f1[idx] <= EYE_R2 ||
						    f2[idx] <= EYE_R2)
							inside = 0;
						if (f3[idx] >= MOUTH_IN &&
						    f3[idx] <= MOUTH_OUT &&
						    v >= MOUTH_V)
							inside = 0;
						hit += inside;
					}
				}
				art[cy][cx] = hit >= CELL * CELL / 2 ? '#' :
									 ' ';
				inked += art[cy][cx] == '#';
			}
			art[cy][CW] = 0;
		}

		printf("RESULT cells=%dx%d inked=%d checksum=%.0f\n",
		       CW, CH, inked, checksum);
		for (cy = 0; cy < CH; cy++)
			printf("%s\n", art[cy]);
	}

	printf("STATUS OK\n");
	CHECK(cudaFree(d_a));
	CHECK(cudaFree(d_b));
	CHECK(cudaFree(d_c));
	free(h_a);
	free(h_b);
	free(h_out);
	free(fields);
	return 0;
}