diff options
Diffstat (limited to 'tools')
| -rw-r--r-- | tools/cuda-ab/victory_mail.cu | 26 |
1 files changed, 20 insertions, 6 deletions
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)); |