// SPDX-License-Identifier: BSD-3-Clause
/*
 * memset_probe.cu - isolate the memcpy_roundtrip memset-phase failure.
 *
 * cudaMemset(8 MB, 0x5A) -> D2H -> byte verify, with a run-length report:
 * the first bad offset, and a coverage map at 4 KB granularity (correct /
 * wrong / untouched-malloc-prefill) so "fill under-delivered" is
 * distinguishable from "readback delivered wrong pages".
 *
 * Build: nvcc -O2 --cudart=static -o memset_probe memset_probe.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)

#define PAGE 4096
#define MemsetByte 0x5A

int main(void)
{
	const size_t bytes = 8u << 20;
	unsigned char *h;
	unsigned char *d;
	size_t i;
	int bad = 0;
	size_t first_bad = (size_t)-1;
	size_t n_pages, p;

	setvbuf(stdout, NULL, _IONBF, 0);
	h = (unsigned char *)malloc(bytes);
	if (!h)
		return 1;
	memset(h, 0xEE, bytes);

	CHECK(cudaMalloc(&d, bytes));
	printf("DPTR d=%p bytes=%zu\n", (void *)d, bytes);

	CHECK(cudaMemset(d, MemsetByte, bytes));
	CHECK(cudaDeviceSynchronize());
	CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost));

	for (i = 0; i < bytes; i++) {
		if (h[i] != (unsigned char)MemsetByte) {
			if (first_bad == (size_t)-1)
				first_bad = i;
			bad++;
		}
	}
	printf("FIRST_BAD off=%zu (0x%zx) bad=%d total=%zu\n",
	       first_bad, first_bad, bad, bytes);
	n_pages = bytes / PAGE;
	for (p = 0; p < n_pages; p++) {
		int good = 1, prefill = 1;

		for (i = 0; i < PAGE; i++) {
			if (h[p * PAGE + i] != (unsigned char)MemsetByte)
				good = 0;
			if (h[p * PAGE + i] != 0xEE)
				prefill = 0;
			if (!good && !prefill)
				break;
		}
		if (!good) {
			printf("PAGE %zu %s b0=%02x b1=%02x b2=%02x b3=%02x b4095=%02x\n",
			       p, prefill ? "PREFILL" : "OTHER",
			       h[p * PAGE], h[p * PAGE + 1],
			       h[p * PAGE + 2], h[p * PAGE + 3],
			       h[p * PAGE + 4095]);
			if (p > 24)
				break;
		}
	}
	printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
	return 0;
}