// SPDX-License-Identifier: BSD-3-Clause
/*
 * memcpy_roundtrip.cu - 8 MB H2D + kernel + D2H roundtrip plus memset check.
 *
 * Two data paths are exercised over an 8 MB (2M float) device buffer:
 *
 *   1. host -> device copy, a kernel (out[i] = in[i] + 0.5), device -> host
 *      copy; the host verifies the full chain element by element and sums the
 *      result as a checksum.
 *   2. cudaMemset of a second 8 MB buffer with a known byte pattern, copied
 *      back and verified byte-for-byte.
 *
 * Output lines are stable and greppable:
 *
 *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
 *   ROUNDTRIP bytes=<n> bad=<mismatches> csum=<sum>
 *   STATUS OK|FAIL
 *
 * Build with a static CUDA runtime (--cudart=static) so the guest only needs
 * the driver's libcuda.so, not a CUDA toolkit installation:
 *
 *   nvcc -O2 --cudart=static -o memcpy_roundtrip memcpy_roundtrip.cu
 *
 * Then, from the guest:  /tmp/memcpy_roundtrip
 */

#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 MemsetByte 0x5A

__global__ void bump(const float *in, float *out, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < n)
		out[i] = in[i] + 0.5f;
}

int main(void)
{
	cudaDeviceProp prop;
	const size_t bytes = 8u << 20; /* 8 MB */
	const int n = (int)(bytes / sizeof(float));
	unsigned char *h_m;
	float *h_in, *h_out;
	float *d_in, *d_out;
	unsigned char *d_m;
	double csum = 0.0;
	int bad = 0, i;
	const int grid = (n + 255) / 256;

	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_in = (float *)malloc(bytes);
	h_out = (float *)malloc(bytes);
	h_m = (unsigned char *)malloc(bytes);
	if (!h_in || !h_out || !h_m)
		return 1;

	for (i = 0; i < n; i++)
		h_in[i] = (float)(i % 1024) * 0.25f;

	CHECK(cudaMalloc(&d_in, bytes));
	CHECK(cudaMalloc(&d_out, bytes));
	CHECK(cudaMalloc(&d_m, bytes));

	/* H2D -> kernel -> D2H roundtrip. */
	CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice));
	bump<<<grid, 256>>>(d_in, d_out, n);
	CHECK(cudaGetLastError());
	CHECK(cudaDeviceSynchronize());
	CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost));

	for (i = 0; i < n; i++) {
		if (h_out[i] != h_in[i] + 0.5f)
			bad++;
		csum += (double)h_out[i];
	}

	/* cudaMemset pattern, copied back and verified byte-for-byte. */
	CHECK(cudaMemset(d_m, MemsetByte, bytes));
	CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost));
	for (i = 0; i < (int)bytes; i++) {
		if (h_m[i] != (unsigned char)MemsetByte)
			bad++;
	}

	printf("ROUNDTRIP bytes=%zu bad=%d csum=%.6f\n", bytes, bad, csum);

	if (bad != 0) {
		printf("STATUS FAIL\n");
		return 1;
	}

	printf("STATUS OK\n");
	CHECK(cudaFree(d_in));
	CHECK(cudaFree(d_out));
	CHECK(cudaFree(d_m));
	free(h_in);
	free(h_out);
	free(h_m);
	return 0;
}