// SPDX-License-Identifier: BSD-3-Clause
/*
 * p2_bench - compute-heavy FMA benchmark with a fixed, verifiable result.
 *
 * 8192 threads each run a 250,000-iteration serial chain of correctly
 * rounded FMAs (__fmaf_rn on device, fmaf on host), producing 2.05e9 FMAs.
 * Both implementations are IEEE-correctly-rounded, so the per-thread final
 * value must match the host exactly (bad=0); the same value is printed for
 * the A/B comparison. A TIMING line (ignored by the harness comparison)
 * reports wall-clock ms and GFLOP/s for the performance demonstration.
 * Device memory: 32 KB output (8 pages).
 */
#include "common.h"
#include <math.h>

#define BENCH_THREADS 8192
#define BENCH_ITERS 250000
#define BENCH_A 0.9999f
#define BENCH_B 1.0e-6f

__global__ void fma_chain(float *out, int iters, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	float x = 0.0f;

	if (i >= n)
		return;
	for (int t = 0; t < iters; t++)
		x = __fmaf_rn(x, BENCH_A, BENCH_B);
	out[i] = x;
}

static float host_chain(int iters)
{
	float x = 0.0f;

	for (int t = 0; t < iters; t++)
		x = fmaf(x, BENCH_A, BENCH_B);
	return x;
}

int main(int argc, char **argv)
{
	float ref = host_chain(BENCH_ITERS);

	if (p2_emit_mode(argc, argv)) {
		printf("RESULT bench threads=%d iters=%d bad=0 value=%.9g\n",
		       BENCH_THREADS, BENCH_ITERS, (double)ref);
		printf("TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU)\n");
		return p2_status(0);
	}

	p2_device_info();

	float *d_out = NULL, *got = (float *)malloc(sizeof(float) * BENCH_THREADS);

	if (!got)
		return 1;
	P2_CHECK(cudaMalloc(&d_out, sizeof(float) * BENCH_THREADS));

	cudaEvent_t t0, t1;
	float ms = 0.0f;

	P2_CHECK(cudaEventCreate(&t0));
	P2_CHECK(cudaEventCreate(&t1));
	P2_CHECK(cudaEventRecord(t0));
	fma_chain<<<BENCH_THREADS / 256, 256>>>(d_out, BENCH_ITERS,
						BENCH_THREADS);
	P2_CHECK(cudaGetLastError());
	P2_CHECK(cudaEventRecord(t1));
	P2_CHECK(cudaDeviceSynchronize());
	P2_CHECK(cudaEventElapsedTime(&ms, t0, t1));
	P2_CHECK(cudaMemcpy(got, d_out, sizeof(float) * BENCH_THREADS,
			    cudaMemcpyDeviceToHost));

	int bad = 0;

	for (int i = 0; i < BENCH_THREADS; i++)
		if (got[i] != ref)
			bad++;
	printf("RESULT bench threads=%d iters=%d bad=%d value=%.9g\n",
	       BENCH_THREADS, BENCH_ITERS, bad, (double)got[0]);
	if (ms > 0.0f) {
		double flops = 2.0 * (double)BENCH_THREADS * (double)BENCH_ITERS;

		printf("TIMING bench ms=%.3f gflops=%.2f\n", (double)ms,
		       flops / ((double)ms * 1.0e6));
	}

	cudaEventDestroy(t0);
	cudaEventDestroy(t1);
	cudaFree(d_out);
	free(got);
	return p2_status(bad);
}