// SPDX-License-Identifier: BSD-3-Clause
/*
 * p2_params - kernel parameter-count and parameter-type coverage.
 *
 * k0..k12 have 0..12 value parameters respectively (plus the output
 * pointer), mixing int, float, double and device-pointer types. Every
 * kernel writes one int slot of a 13-int device buffer; the CPU reference
 * mirrors each expression exactly and the buffer must match slot for slot.
 *
 * Device memory: out 52 B, p 8 B, pf 4 B (1 page).
 */
#include "common.h"

#define NK 13

__global__ void k0(int *out)
{
	out[0] = 7;
}

__global__ void k1(int *out, int a)
{
	out[1] = a + 1;
}

__global__ void k2(int *out, int a, int b)
{
	out[2] = a + b;
}

__global__ void k3(int *out, int a, int b, float f1)
{
	out[3] = a + b + (int)f1;
}

__global__ void k4(int *out, int a, int b, float f1, double c1)
{
	out[4] = a + b + (int)f1 + (int)c1;
}

__global__ void k5(int *out, int a, int b, float f1, double c1, int d)
{
	out[5] = a + b + (int)f1 + (int)c1 + d;
}

__global__ void k6(int *out, int a, int b, float f1, double c1, int d,
		   const int *p)
{
	out[6] = a + b + (int)f1 + (int)c1 + d + p[0];
}

__global__ void k7(int *out, int a, int b, float f1, double c1, int d,
		   const int *p, const int *q)
{
	out[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0];
}

__global__ void k8(int *out, int a, int b, float f1, double c1, int d,
		   const int *p, const int *q, int n)
{
	out[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n;
}

__global__ void k9(int *out, int a, int b, float f1, double c1, int d,
		   const int *p, const int *q, int n, float f2)
{
	out[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2;
}

__global__ void k10(int *out, int a, int b, float f1, double c1, int d,
		    const int *p, const int *q, int n, float f2,
		    double c2)
{
	out[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
		  (int)f2 + (int)c2;
}

__global__ void k11(int *out, int a, int b, float f1, double c1, int d,
		    const int *p, const int *q, int n, float f2,
		    double c2, float s)
{
	out[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
		  (int)f2 + (int)c2 + (int)s;
}

__global__ void k12(int *out, int a, int b, float f1, double c1, int d,
		    const int *p, const int *q, int n, float f2,
		    double c2, float s, int e)
{
	out[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
		  (int)f2 + (int)c2 + (int)s + e;
}

/* CPU mirror of the 13 kernel expressions; keeps expected values exact. */
static void expected(int *exp, int a, int b, float f1, double c1, int d,
		     const int *p, const int *q, int n, float f2, double c2,
		     float s, int e)
{
	exp[0] = 7;
	exp[1] = a + 1;
	exp[2] = a + b;
	exp[3] = a + b + (int)f1;
	exp[4] = a + b + (int)f1 + (int)c1;
	exp[5] = a + b + (int)f1 + (int)c1 + d;
	exp[6] = a + b + (int)f1 + (int)c1 + d + p[0];
	exp[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0];
	exp[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n;
	exp[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2;
	exp[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
		  (int)f2 + (int)c2;
	exp[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
		  (int)f2 + (int)c2 + (int)s;
	exp[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
		  (int)f2 + (int)c2 + (int)s + e;
}

static int checksum(const int *v)
{
	int sum = 0;

	for (int i = 0; i < NK; i++)
		sum += v[i];
	return sum;
}

int main(int argc, char **argv)
{
	int a = 3, b = 5, d = 11, n = 2, e = 4;
	float f1 = 6.5f, f2 = 8.5f, s = 2.5f;
	double c1 = 2.5, c2 = 3.5;
	int hp[2] = { 5, 6 }, hq[2] = { 6, 7 };
	int exp[NK];

	expected(exp, a, b, f1, c1, d, hp, hq, n, f2, c2, s, e);

	if (p2_emit_mode(argc, argv)) {
		printf("RESULT params kernels=%d nparam_max=12 bad=0 checksum=%d\n",
		       NK, checksum(exp));
		return p2_status(0);
	}

	p2_device_info();

	int *d_out = NULL, *got = (int *)malloc(sizeof(int) * NK);
	int *dp = NULL, *dq = NULL;

	if (!got)
		return 1;
	P2_CHECK(cudaMalloc(&d_out, sizeof(int) * NK));
	P2_CHECK(cudaMalloc(&dp, sizeof(int) * 2));
	P2_CHECK(cudaMalloc(&dq, sizeof(int) * 2));
	P2_CHECK(cudaMemcpy(dp, hp, sizeof(int) * 2, cudaMemcpyHostToDevice));
	P2_CHECK(cudaMemcpy(dq, hq, sizeof(int) * 2, cudaMemcpyHostToDevice));
	P2_CHECK(cudaMemset(d_out, 0, sizeof(int) * NK));

	k0<<<1, 1>>>(d_out);
	k1<<<1, 1>>>(d_out, a);
	k2<<<1, 1>>>(d_out, a, b);
	k3<<<1, 1>>>(d_out, a, b, f1);
	k4<<<1, 1>>>(d_out, a, b, f1, c1);
	k5<<<1, 1>>>(d_out, a, b, f1, c1, d);
	k6<<<1, 1>>>(d_out, a, b, f1, c1, d, dp);
	k7<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq);
	k8<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n);
	k9<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2);
	k10<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2);
	k11<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s);
	k12<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s, e);
	P2_CHECK(cudaGetLastError());
	P2_CHECK(cudaDeviceSynchronize());
	P2_CHECK(cudaMemcpy(got, d_out, sizeof(int) * NK,
			    cudaMemcpyDeviceToHost));

	int bad = 0;

	for (int i = 0; i < NK; i++)
		if (got[i] != exp[i])
			bad++;
	printf("RESULT params kernels=%d nparam_max=12 bad=%d checksum=%d\n",
	       NK, bad, checksum(got));

	cudaFree(d_out); cudaFree(dp); cudaFree(dq);
	free(got);
	return p2_status(bad);
}