// SPDX-License-Identifier: BSD-3-Clause
/*
 * memset2d_probe.cu - per-phase diagnostics for victory_memset2d.
 *
 * victory_memset2d reports one aggregate bad count (~12.5M) that mixes all
 * three phases. This probe runs the SAME three phases and reports each
 * phase's bad count, the first mismatching byte (offset, got, want), and a
 * per-phase 1MB-chunk histogram so the failing phase and the failing region
 * are attributable in one run.
 *
 * Phases:
 *   P1 memset   : cudaMemset(d_img, 0xA5, pitch*h); D2H 2D readback
 *   P2 roundtrip: H2D 2D pattern; D2H 2D readback
 *   P3 memset2d : cudaMemset2D quadrant 0x3C; D2H 2D readback
 *
 * Output lines:
 *   DEVICE ...
 *   PHASE p=1 bad=<n> first_off=<i> got=<g> want=<w>
 *   PHIST p=1 mb=<m> ok=<n> bad=<n>
 *   MEMSET2D2 bad_total=<n>
 *   STATUS OK|FAIL
 */

#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cstring>
#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)

#define IMG_W 4096
#define IMG_H 1024
#define FILL_BYTE 0xA5
#define QUAD_BYTE 0x3C

static int phase_check(const char *tag, const unsigned char *back,
		       const unsigned char *want, size_t n, int p)
{
	size_t i, first = (size_t)-1;
	int bad = 0, got = 0, wantv = 0;
	size_t mb;

	for (i = 0; i < n; i++) {
		if (back[i] != want[i]) {
			if (first == (size_t)-1) {
				first = i;
				got = back[i];
				wantv = want[i];
			}
			bad++;
		}
	}
	printf("PHASE p=%d bad=%d first_off=%zu got=0x%02x want=0x%02x\n",
	       p, bad, first, got, wantv);
	if (bad) {
		size_t nb = (n + (1u << 20) - 1) >> 20;

		for (mb = 0; mb < nb && mb < 8; mb++) {
			size_t lo = mb << 20, hi = lo + (1u << 20);
			size_t j;
			int bok = 0, bbad = 0;

			if (hi > n)
				hi = n;
			for (j = lo; j < hi; j++) {
				if (back[j] == want[j])
					bok++;
				else
					bbad++;
			}
			printf("PHIST p=%d mb=%zu ok=%d bad=%d\n",
			       p, mb, bok, bbad);
		}
	}
	(void)tag;
	return bad;
}

int main(void)
{
	cudaDeviceProp prop;
	const size_t w = IMG_W, h = IMG_H;
	unsigned char *h_img, *h_back, *h_want;
	unsigned char *d_img;
	size_t pitch = 0;
	int bad = 0, row, col;

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

	h_img = (unsigned char *)malloc(w * h);
	h_back = (unsigned char *)malloc(w * h);
	h_want = (unsigned char *)malloc(w * h);
	if (!h_img || !h_back || !h_want)
		return 1;

	CHECK(cudaMallocPitch(&d_img, &pitch, w, h));
	printf("PITCH pitch=%zu\n", pitch);
	printf("DPTRS h_img=%p h_back=%p h_want=%p d_img=%p\n", (void *)h_img,
	       (void *)h_back, (void *)h_want, (void *)d_img);

	/* P1: full cudaMemset over the pitched allocation, 2D readback. */
	CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h));
	memset(h_back, 0, w * h);
	CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
			   cudaMemcpyDeviceToHost));
	memset(h_want, FILL_BYTE, w * h);
	bad += phase_check("P1", h_back, h_want, w * h, 1);

	/* P2: pitched H2D pattern -> D2H round trip. */
	for (row = 0; row < (int)h; row++)
		for (col = 0; col < (int)w; col++)
			h_img[row * w + col] = (unsigned char)((row + col) & 0xFF);
	CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h,
			   cudaMemcpyHostToDevice));
	memset(h_back, 0, w * h);
	CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
			   cudaMemcpyDeviceToHost));
	bad += phase_check("P2", h_back, h_img, w * h, 2);

	/* P3: cudaMemset2D of the upper-left quadrant only. */
	CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2));
	for (row = 0; row < (int)h; row++)
		for (col = 0; col < (int)w; col++) {
			unsigned char expect = h_img[row * w + col];

			if (row < (int)(h / 2) && col < (int)(w / 2))
				expect = QUAD_BYTE;
			h_want[row * w + col] = expect;
		}
	memset(h_back, 0, w * h);
	CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
			   cudaMemcpyDeviceToHost));
	bad += phase_check("P3", h_back, h_want, w * h, 3);

	printf("MEMSET2D2 bad_total=%d\n", bad);
	if (bad != 0) {
		printf("STATUS FAIL\n");
		return 1;
	}
	printf("STATUS OK\n");
	return 0;
}