1 files changed, 123 insertions, 0 deletions
diff --git a/tools/cuda-ab/memcpy_roundtrip.cu b/tools/cuda-ab/memcpy_roundtrip.cu
new file mode 100644
index 0000000..679181f
--- /dev/null
+++ b/tools/cuda-ab/memcpy_roundtrip.cu
@@ -0,0 +1,123 @@
+// 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;
+}