From d091097c03d2db3f8100574a6f8b1a7a4b5b2a01 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Wed, 7 Oct 2026 15:49:31 -0400 Subject: cuda-ab: victory_mail native-path sizing + param-based mail message Two changes for the native GR path: vadd buffers shrink to 1<<15 elements (128 KB each) so all data fits the engine-visible 2 MB window, and mail_kernel receives the message as a kernel-parameter buffer instead of a __device__ constant (the c[0x4] module-data base is driver-populated and is zero in the forwarded launch). The guest binary must be rebuilt in tools/cuda-ab/run (ab.sh push_bin only builds when the binary is absent). --- tools/cuda-ab/victory_mail.cu | 26 ++++++++++++++++++++------ 1 file changed, 20 insertions(+), 6 deletions(-) (limited to 'tools') diff --git a/tools/cuda-ab/victory_mail.cu b/tools/cuda-ab/victory_mail.cu index e8136c3..e9f2ed2 100644 --- a/tools/cuda-ab/victory_mail.cu +++ b/tools/cuda-ab/victory_mail.cu @@ -51,16 +51,24 @@ } \ } while (0) -/* The message lives in device code: the GPU both provides and writes it. */ -__device__ const char kMail[] = MAIL; +/* The message lives in the host binary. */ +static const char kMail[] = MAIL; -/* Each thread writes exactly one character of the message. */ -__global__ void mail_kernel(char *out, int len) +/* + * 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] = kMail[i]; + out[i] = msg[i]; } __global__ void vadd(const float *a, const float *b, float *c, int n) @@ -134,7 +142,13 @@ int main(void) /* The GPU writes the mail. */ CHECK(cudaMalloc(&d_mail, len)); dump_va("d_mail_alloc", d_mail); - mail_kernel<<<1, 64>>>(d_mail, len); + { + 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)); -- cgit v1.2.3