// SPDX-License-Identifier: BSD-3-Clause
/*
 * inplace_scale.cu - minimal in-place scale repro for the streams2
 * "scale wrote 1 page" class.
 *
 * One H2D, one in-place scale launch (c[i] *= 2 over 128 KB = 32 pages),
 * one D2H, per-page diff. Run it repeatedly: if the launch writes only
 * part of the buffer the per-page report names how many 4 KB blocks
 * changed and the host log carries the launch's retire wait.
 *
 * Build: nvcc -O2 --cudart=static -o inplace_scale inplace_scale.cu
 */

#include <cstdio>
#include <cstdlib>
#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 scale(float *c, float s, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

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

int main(void)
{
	const int n = 1 << 15;
	size_t bytes = (size_t)n * sizeof(float);
	float *h, *d;
	int i, bad = 0, first = -1, last = -1;
	int nb = (n + 1023) / 1024, blk_bad = 0;

	setvbuf(stdout, NULL, _IONBF, 0);
	h = (float *)malloc(bytes);
	if (!h)
		return 1;
	for (i = 0; i < n; i++)
		h[i] = (float)(i % 1000) * 0.5f;

	CHECK(cudaMalloc(&d, bytes));
	CHECK(cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice));
	printf("PTR d=%p\n", (void *)d);

	scale<<<(n + 255) / 256, 256>>>(d, 2.0f, n);
	CHECK(cudaGetLastError());
	CHECK(cudaDeviceSynchronize());

	CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost));
	for (i = 0; i < n; i++) {
		float exp = (float)(i % 1000) * 0.5f * 2.0f;

		if (h[i] != exp) {
			bad++;
			if (first < 0)
				first = i;
			last = i;
		}
	}
	for (i = 0; i < nb; i++) {
		int j, hit = 0;

		for (j = i * 1024; j < (i + 1) * 1024 && j < n; j++)
			if (h[j] != (float)(j % 1000) * 0.5f * 2.0f) {
				hit = 1;
				break;
			}
		if (hit)
			blk_bad++;
	}
	printf("SCALE n=%d bad=%d first=%d last=%d blocks=%d/%d\n",
	       n, bad, first, last, blk_bad, nb);
	printf("STATUS %s\n", bad ? "FAIL" : "OK");
	return bad ? 1 : 0;
}