// SPDX-License-Identifier: BSD-3-Clause
/*
 * p2_streams - three created streams, two kernels per stream (scale then
 * add), independent 1024-float buffers, per-stream events.
 *
 * Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32).
 * After all streams complete, each event must query as complete.
 * Device memory: 3 x 4 KB (3 pages).
 */
#include "common.h"

#define SN 1024
#define NS 3

__global__ void scale(float *x, float s, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < n)
		x[i] = x[i] * s;
}

__global__ void addv(float *x, float v, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < n)
		x[i] = x[i] + v;
}

static void fill(float *h, int n)
{
	for (int i = 0; i < n; i++)
		h[i] = (float)(i % 100);
}

static float expected1(int i)
{
	return ((float)(i % 100) * 2.0f) + 1.0f;
}

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

	if (!h || !got)
		return 1;
	fill(h, SN);

	if (p2_emit_mode(argc, argv)) {
		printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n",
		       NS, NS * 2, SN);
		printf("RESULT streams events=%d bad=0\n", NS);
		free(h); free(got);
		return p2_status(0);
	}

	p2_device_info();

	cudaStream_t st[NS];
	cudaEvent_t ev[NS];
	float *d_x[NS];

	for (int k = 0; k < NS; k++) {
		P2_CHECK(cudaStreamCreate(&st[k]));
		P2_CHECK(cudaEventCreate(&ev[k]));
		P2_CHECK(cudaMalloc(&d_x[k], sizeof(float) * SN));
		P2_CHECK(cudaMemcpyAsync(d_x[k], h, sizeof(float) * SN,
					 cudaMemcpyHostToDevice, st[k]));
		scale<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 2.0f, SN);
		addv<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 1.0f, SN);
		P2_CHECK(cudaGetLastError());
		P2_CHECK(cudaEventRecord(ev[k], st[k]));
	}

	int ibad = 0;

	for (int k = 0; k < NS; k++) {
		cudaError_t q;
		int sbad = 0;

		P2_CHECK(cudaStreamSynchronize(st[k]));
		P2_CHECK(cudaMemcpy(got, d_x[k], sizeof(float) * SN,
				    cudaMemcpyDeviceToHost));
		for (int i = 0; i < SN; i++)
			if (got[i] != expected1(i))
				sbad++;
		q = cudaEventQuery(ev[k]);
		if (q != cudaSuccess)
			sbad++;
		printf("RESULT streams stream=%d query_rc=%d bad=%d\n", k, (int)q,
		       sbad);
		ibad += sbad;
	}
	printf("RESULT streams nstreams=%d launches=%d n=%d bad=%d\n",
	       NS, NS * 2, SN, ibad);
	printf("RESULT streams events=%d bad=%d\n", NS, ibad == 0 ? 0 : 1);
	bad = ibad;

	for (int k = 0; k < NS; k++) {
		cudaEventDestroy(ev[k]);
		cudaStreamDestroy(st[k]);
		cudaFree(d_x[k]);
	}
	free(h); free(got);
	return p2_status(bad);
}