From c8ccdd772c7650f19d14f143f34b700813575746 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Thu, 8 Oct 2026 04:15:24 -0400 Subject: vgsp+guest-diag: kcore-truth pin resolver (pool/pushbuffer/channel ring+userd) + GR ring userd-derived candidates guest_pin_stream.py now resolves translations entirely from guest truth via /proc/kcore + a manual page-table walk (no QMP races): the CE semaphore pool and pushbuffer memory from the UVM rm_mem structures, and each channel's GPFIFO ring + GPPUT from the tracking semaphore's embedding uvm_channel_t (channel_info located by its gpFifoGpuVa value). This removes the watcher pipeline's flakiness on the cudaMalloc stall class (missing pool/pushbuffer pins). gsp_sema_service.c: svc_gr_ring_autopin tries the guest-registered USERD memdesc base as a ring candidate (the fixed FB convention only matches when the guest's VRAM allocator happens to reproduce the old layout), plus an optional (gsp_sema_dbg) one-shot FB scan + raw-word diagnostics when every candidate fails. tools/cuda-ab/p2: self-checking arbitrary-kernel acceptance suite (10 programs: sequential launches, shapes, params, smem/GEMM, streams, malloc churn, printf+error, FMA benchmark, mixed, managed probe) with host oracle expectations and a guest runner. --- tools/cuda-ab/p2/src/p2_seq5.cu | 127 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 127 insertions(+) create mode 100644 tools/cuda-ab/p2/src/p2_seq5.cu (limited to 'tools/cuda-ab/p2/src/p2_seq5.cu') diff --git a/tools/cuda-ab/p2/src/p2_seq5.cu b/tools/cuda-ab/p2/src/p2_seq5.cu new file mode 100644 index 0000000..8a778f4 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_seq5.cu @@ -0,0 +1,127 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_seq5 - five sequential kernel launches in one process on the same + * buffers, verifying the device state after every launch. + * + * step 1: c = a + b + * step 2: c = c + a + * step 3: c = c + b + * step 4: c = c * 2 + * step 5: c = c + 1 + * + * Inputs are small integers so every operation is exact in fp32; the CPU + * reference reproduces the same chain and must match element for element. + * Device memory: 3 x 4 KB pointers (3 pages total). + */ +#include "common.h" + +#define N 1024 + +__global__ void kadd(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]; +} + +__global__ void kmul2(float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * 2.0f; +} + +__global__ void kaddv(float *c, float v, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] + v; +} + +static void make_inputs(float *a, float *b) +{ + for (int i = 0; i < N; i++) { + a[i] = (float)(i % 7); + b[i] = (float)(i % 5); + } +} + +static void step_cpu(float *c, const float *a, const float *b, int step) +{ + for (int i = 0; i < N; i++) { + switch (step) { + case 1: c[i] = a[i] + b[i]; break; + case 2: c[i] = c[i] + a[i]; break; + case 3: c[i] = c[i] + b[i]; break; + case 4: c[i] = c[i] * 2.0f; break; + case 5: c[i] = c[i] + 1.0f; break; + } + } +} + +int main(int argc, char **argv) +{ + float *h_a, *h_b, *h_c, *h_exp; + int bad = 0; + + if (p2_emit_mode(argc, argv)) { + h_a = (float *)malloc(sizeof(float) * N); + h_b = (float *)malloc(sizeof(float) * N); + h_exp = (float *)malloc(sizeof(float) * N); + make_inputs(h_a, h_b); + for (int step = 1; step <= 5; step++) { + step_cpu(h_exp, h_a, h_b, step); + printf("RESULT seq step=%d n=%d bad=0\n", step, N); + } + printf("RESULT seq launches=5 n=%d bad=0\n", N); + free(h_a); free(h_b); free(h_exp); + return p2_status(0); + } + + p2_device_info(); + h_a = (float *)malloc(sizeof(float) * N); + h_b = (float *)malloc(sizeof(float) * N); + h_c = (float *)malloc(sizeof(float) * N); + h_exp = (float *)malloc(sizeof(float) * N); + if (!h_a || !h_b || !h_c || !h_exp) + return 1; + make_inputs(h_a, h_b); + + float *d_a = NULL, *d_b = NULL, *d_c = NULL; + + P2_CHECK(cudaMalloc(&d_a, sizeof(float) * N)); + P2_CHECK(cudaMalloc(&d_b, sizeof(float) * N)); + P2_CHECK(cudaMalloc(&d_c, sizeof(float) * N)); + P2_CHECK(cudaMemcpy(d_a, h_a, sizeof(float) * N, cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemcpy(d_b, h_b, sizeof(float) * N, cudaMemcpyHostToDevice)); + + for (int step = 1; step <= 5; step++) { + int sbad = 0; + + switch (step) { + case 1: kadd<<<(N + 255) / 256, 256>>>(d_a, d_b, d_c, N); break; + case 2: kadd<<<(N + 255) / 256, 256>>>(d_c, d_a, d_c, N); break; + case 3: kadd<<<(N + 255) / 256, 256>>>(d_c, d_b, d_c, N); break; + case 4: kmul2<<<(N + 255) / 256, 256>>>(d_c, N); break; + case 5: kaddv<<<(N + 255) / 256, 256>>>(d_c, 1.0f, N); break; + } + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(h_c, d_c, sizeof(float) * N, + cudaMemcpyDeviceToHost)); + step_cpu(h_exp, h_a, h_b, step); + for (int i = 0; i < N; i++) + if (h_c[i] != h_exp[i]) + sbad++; + printf("RESULT seq step=%d n=%d bad=%d\n", step, N, sbad); + bad += sbad; + } + printf("RESULT seq launches=5 n=%d bad=%d\n", N, bad); + + cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); + free(h_a); free(h_b); free(h_c); free(h_exp); + return p2_status(bad); +} -- cgit v1.2.3