path: root/tools/cuda-ab/victory_mail.cu

// SPDX-License-Identifier: BSD-3-Clause
/*
 * victory_mail.cu - the Open-IOV victory program.
 *
 * Run from the guest. Two kernels execute on the physical GPU:
 *
 *   mail_kernel  - the GPU itself writes every character of the message into
 *                  device memory; the host only copies the bytes back.
 *   vadd         - the classic vectorAdd acceptance check, verified element
 *                  by element against the host-computed result.
 *
 * The host prints the message only after it has been authored by GPU
 * execution and copied back byte-for-byte, so the line below is proof that a
 * kernel really ran on the device:
 *
 *   the little gpu man says you've got mail
 *
 * Output lines are stable and greppable:
 *
 *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
 *   MAIL <message>
 *   VADD n=<count> bad=<mismatches> checksum=<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 victory_mail victory_mail.cu
 *
 * Then, from the guest:  /tmp/victory_mail
 */

#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cstdint>
#include <fcntl.h>
#include <unistd.h>
#include <cuda_runtime.h>

#define MAIL "the little gpu man says you've got mail"

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

/* The message lives in the host binary. */
static const char kMail[] = MAIL;

/*
 * m115 (native GR path): the message is passed as a kernel parameter buffer,
 * like vadd's inputs, instead of a __device__ constant. Reading a
 * __device__ array makes the compiler load its base from const bank 4
 * (c[0x4][0x0]), which the guest driver never populated in the forwarded
 * launch (zero in guest memory); the parameter path is populated by the
 * driver's own memcpy and works on both the native and bridge paths.
 * Each thread writes exactly one character of the message.
 */
__global__ void mail_kernel(char *out, const char *msg, int len)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < len)
		out[i] = msg[i];
}

__global__ void vadd(const float *a, const float *b, float *c, int n)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	if (i < n)
		c[i] = a[i] + b[i];
}

/* Diagnostic: where does a CUDA pointer live for the CPU?
 * Prints the /proc/self/maps line covering it and the pagemap
 * present/PFN entry, so the read-back mechanism of the DtoH copy is on
 * record (CPU mapping vs device-only). Never affects control flow. */
static void dump_va(const char *tag, const void *p)
{
	unsigned long long want = (unsigned long long)(uintptr_t)p;
	char line[512];
	FILE *m = fopen("/proc/self/maps", "r");

	printf("DMAP %s va=0x%llx\n", tag, want);
	if (m) {
		while (fgets(line, sizeof(line), m)) {
			unsigned long long lo, hi;

			if (sscanf(line, "%llx-%llx", &lo, &hi) == 2 &&
			    want >= lo && want < hi) {
				printf("DMAPLINE %s", line);
				break;
			}
		}
		fclose(m);
	}
	int pm = open("/proc/self/pagemap", O_RDONLY);
	if (pm >= 0) {
		unsigned long long e = 0;
		off_t off = (off_t)((want / 4096) * 8);

		if (pread(pm, &e, 8, off) == 8)
			printf("DPRESENT %s present=%d pfn=0x%llx\n", tag,
			       (int)((e >> 63) & 1),
			       (unsigned long long)(e & ((1ULL << 55) - 1)));
		else
			printf("DPRESENT %s read failed\n", tag);
		close(pm);
	} else {
		printf("DPRESENT %s pagemap open failed\n", tag);
	}
}

int main(void)
{
	cudaDeviceProp prop;
	const int len = (int)strlen(MAIL);
	const int n = 1 << 15; /* sized for the native GR path: the engine-visible data window is 2 MB */
	size_t bytes = (size_t)n * sizeof(float);
	char *d_mail, h_mail[128];
	float *h_a, *h_b, *h_c;
	float *d_a, *d_b, *d_c;
	double sum = 0.0;
	int bad = 0, i;

	/* Unbuffered: progress must be visible even if a later call blocks. */
	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);

	/* The GPU writes the mail. */
	CHECK(cudaMalloc(&d_mail, len));
	dump_va("d_mail_alloc", d_mail);
	{
		char *d_msg = NULL;

		CHECK(cudaMalloc(&d_msg, len));
		CHECK(cudaMemcpy(d_msg, kMail, len, cudaMemcpyHostToDevice));
		mail_kernel<<<1, 64>>>(d_mail, d_msg, len);
	}
	CHECK(cudaGetLastError());
	CHECK(cudaDeviceSynchronize());
	memset(h_mail, 0, sizeof(h_mail));
	CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost));
	dump_va("d_mail_after_dtoh", d_mail);
	if (memcmp(h_mail, MAIL, len) != 0) {
		printf("MAILDUMP ");
		for (i = 0; i < len; i++)
			printf("%02x", (unsigned char)h_mail[i]);
		printf("\n");
		printf("MAIL bad=<kernel output did not match>\n");
		printf("STATUS FAIL\n");
		return 1;
	}
	printf("MAIL %s\n", h_mail);
	{
		size_t fm = 0, tm = 0;
		cudaError_t me = cudaMemGetInfo(&fm, &tm);

		printf("MEMINFO rc=%d free=%llu total=%llu\n", (int)me,
		       (unsigned long long)fm, (unsigned long long)tm);
	}

	/* The classic acceptance check: vectorAdd. */
	h_a = (float *)malloc(bytes);
	h_b = (float *)malloc(bytes);
	h_c = (float *)malloc(bytes);
	if (!h_a || !h_b || !h_c)
		return 1;

	for (i = 0; i < n; i++) {
		h_a[i] = (float)(i % 1000) * 0.5f;
		h_b[i] = (float)(i % 997) * 0.25f;
	}

	{
		cudaError_t e_a = cudaMalloc(&d_a, bytes);
		cudaError_t e_b = e_a == cudaSuccess ?
			cudaMalloc(&d_b, bytes) : e_a;
		cudaError_t e_c = e_b == cudaSuccess ?
			cudaMalloc(&d_c, bytes) : e_b;

		printf("MALLOC bytes=%zu d_a=%d d_b=%d d_c=%d\n", bytes,
		       (int)e_a, (int)e_b, (int)e_c);
		if (e_a != cudaSuccess) {
			size_t fm = 0, tm = 0;

			cudaMemGetInfo(&fm, &tm);
			printf("MEMINFO2 free=%llu total=%llu\n",
			       (unsigned long long)fm,
			       (unsigned long long)tm);
		}
		CHECK(e_a);
		CHECK(e_b);
		CHECK(e_c);
	}
	CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
	CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

	vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
	CHECK(cudaGetLastError());
	CHECK(cudaDeviceSynchronize());
	CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));

	for (i = 0; i < n; i++) {
		if (h_c[i] != h_a[i] + h_b[i])
			bad++;
		sum += h_c[i];
	}

	printf("VADD n=%d bad=%d checksum=%.6f\n", n, bad, sum);

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

	printf("the little gpu man says you've got mail\n");
	printf("STATUS OK\n");

	CHECK(cudaFree(d_mail));
	CHECK(cudaFree(d_a));
	CHECK(cudaFree(d_b));
	CHECK(cudaFree(d_c));
	free(h_a);
	free(h_b);
	free(h_c);
	return 0;
}