diff options
Diffstat (limited to 'tools')
57 files changed, 6429 insertions, 120 deletions
diff --git a/tools/cuda-ab/evhold.cu b/tools/cuda-ab/evhold.cu new file mode 100644 index 0000000..3817ff3 --- /dev/null +++ b/tools/cuda-ab/evhold.cu @@ -0,0 +1,56 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * evhold.cu - deterministic event-holder / event-wait probe for the guest + * event-delivery class (r577). Opens a CUDA context (so the driver + * registers its RM events), runs one small op, records and synchronizes a + * CUDA event (which exercises the client's wait path), then holds the + * context open for N seconds so host-side experiments (rmpost rounds with + * the gsp_rm_ev_* knobs, ev_post kprobe) have a live event table to post + * into. + * + * Usage: evhold [hold_seconds] + * Output lines: + * EVHOLD ctx ok + * EVHOLD event sync ok + * EVHOLD hold done + */ +#include <cstdio> +#include <cstdlib> +#include <unistd.h> +#include <cuda_runtime.h> + +static void check(const char *what, cudaError_t e) +{ + if (e != cudaSuccess) { + printf("EVHOLD FAIL %s: %s\n", what, cudaGetErrorString(e)); + fflush(stdout); + exit(1); + } +} + +int main(int argc, char **argv) +{ + int hold = (argc > 1) ? atoi(argv[1]) : 120; + float *d = NULL; + cudaEvent_t ev; + + check("freectx", cudaFree(0)); + printf("EVHOLD ctx ok\n"); + fflush(stdout); + + check("malloc", cudaMalloc(&d, 4096)); + check("memset", cudaMemset(d, 0, 4096)); + check("evcreate", cudaEventCreate(&ev)); + check("evrecord", cudaEventRecord(ev, 0)); + check("evsync", cudaEventSynchronize(ev)); + printf("EVHOLD event sync ok\n"); + fflush(stdout); + + for (int i = 0; i < hold; i++) + sleep(1); + + cudaEventDestroy(ev); + cudaFree(d); + printf("EVHOLD hold done\n"); + return 0; +} diff --git a/tools/cuda-ab/frame_producer.cu b/tools/cuda-ab/frame_producer.cu new file mode 100644 index 0000000..609114c --- /dev/null +++ b/tools/cuda-ab/frame_producer.cu @@ -0,0 +1,172 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * frame_producer.cu - render an animated frame on the guest GPU and publish + * it into the ivshmem (kvmfr) BAR so the host can see a real GPU-rendered + * frame from the guest. + * + * Path: CUDA kernel -> device buffer -> host buffer -> ivshmem BAR + * (/sys/bus/pci/devices/0000:00:03.0/resource2, the 128 MiB kvmfr + * device QEMU backs with /dev/shm/kvmfr-<uuid>). + * + * Layout in the BAR (offset 0x100000, clear of the kvmfr protocol header): + * struct frame_hdr { magic, version, width, height, stride, format, seq } + * then width*height*4 bytes of RGBA8. + * + * Usage: sudo ./frame_producer [frames] [width] [height] [fps] + * Prints: FRAME seq=<n> ms=<render+copy ms> (one line per frame, first 5 and + * every 30th) and FRAME_DONE frames=<n>. + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#include <ctime> +#include <fcntl.h> +#include <sys/mman.h> +#include <unistd.h> +#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__); \ + return 1; \ + } \ + } while (0) + +#define FRAME_OFFSET 0x100000UL +#define HDR_MAGIC 0x4b564d46u /* "FMVK" little-endian-ish */ + +struct frame_hdr { + uint32_t magic; + uint32_t version; + uint32_t width; + uint32_t height; + uint32_t stride; + uint32_t format; /* 0 = RGBA8 */ + uint32_t seq; + uint32_t reserved; + uint64_t timestamp_ns; +}; + +__global__ void render_julia(uchar4 *fb, int w, int h, float t) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= w || y >= h) + return; + + float fx = (x - w * 0.5f) / (h * 0.45f); + float fy = (y - h * 0.5f) / (h * 0.45f); + float cr = 0.7885f * cosf(t); + float ci = 0.7885f * sinf(t); + float zr = fx, zi = fy; + int it = 0; + const int maxit = 96; + + while (zr * zr + zi * zi < 4.0f && it < maxit) { + float nzr = zr * zr - zi * zi + cr; + float nzi = 2.0f * zr * zi + ci; + + zr = nzr; + zi = nzi; + it++; + } + { + float v = (float)it / (float)maxit; + unsigned char r = (unsigned char)(255.0f * v); + unsigned char g = (unsigned char)(255.0f * (1.0f - v) * v * 4.0f); + unsigned char b = (unsigned char)(255.0f * (1.0f - v)); + unsigned char a = 255; + + fb[y * w + x] = make_uchar4(r, g, b, a); + } +} + +static uint64_t now_ns(void) +{ + struct timespec ts; + + clock_gettime(CLOCK_MONOTONIC, &ts); + return (uint64_t)ts.tv_sec * 1000000000ull + ts.tv_nsec; +} + +int main(int argc, char **argv) +{ + int frames = argc > 1 ? atoi(argv[1]) : 60; + int w = argc > 2 ? atoi(argv[2]) : 800; + int h = argc > 3 ? atoi(argv[3]) : 600; + int fps = argc > 4 ? atoi(argv[4]) : 10; + const char *bar_path = + "/sys/bus/pci/devices/0000:00:03.0/resource2"; + size_t fb_bytes = (size_t)w * h * 4; + uchar4 *d_fb = NULL, *h_fb = NULL; + unsigned char *bar; + struct frame_hdr *hdr; + volatile uint32_t *seqp; + int fd, i; + + setvbuf(stdout, NULL, _IONBF, 0); + + fd = open(bar_path, O_RDWR | O_SYNC); + if (fd < 0) { + printf("OPEN_FAIL %s\n", bar_path); + return 1; + } + bar = (unsigned char *)mmap(NULL, 128UL << 20, PROT_READ | PROT_WRITE, + MAP_SHARED, fd, 0); + if (bar == MAP_FAILED) { + printf("MMAP_FAIL\n"); + return 1; + } + hdr = (struct frame_hdr *)(bar + FRAME_OFFSET); + seqp = &hdr->seq; + + CHECK(cudaMalloc(&d_fb, fb_bytes)); + h_fb = (uchar4 *)malloc(fb_bytes); + if (!h_fb) { + printf("HOST_ALLOC_FAIL\n"); + return 1; + } + + hdr->magic = HDR_MAGIC; + hdr->version = 1; + hdr->width = (uint32_t)w; + hdr->height = (uint32_t)h; + hdr->stride = (uint32_t)w * 4; + hdr->format = 0; + + for (i = 0; i < frames; i++) { + uint64_t t0 = now_ns(); + float phase = (float)i * 0.15f; + dim3 block(16, 16); + dim3 grid((w + 15) / 16, (h + 15) / 16); + + render_julia<<<grid, block>>>(d_fb, w, h, phase); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_fb, d_fb, fb_bytes, + cudaMemcpyDeviceToHost)); + memcpy(bar + FRAME_OFFSET + sizeof(struct frame_hdr), h_fb, + fb_bytes); + hdr->timestamp_ns = now_ns(); + __sync_synchronize(); + *seqp = (uint32_t)(i + 1); + __sync_synchronize(); + if (i < 5 || (i % 30) == 29) + printf("FRAME seq=%d ms=%.1f\n", i + 1, + (double)(now_ns() - t0) / 1e6); + if (fps > 0) + usleep((useconds_t)(1000000 / fps)); + } + printf("FRAME_DONE frames=%d %dx%d\n", frames, w, h); + CHECK(cudaFree(d_fb)); + free(h_fb); + munmap(bar, 128UL << 20); + close(fd); + return 0; +} diff --git a/tools/cuda-ab/inplace_scale.cu b/tools/cuda-ab/inplace_scale.cu new file mode 100644 index 0000000..e83a039 --- /dev/null +++ b/tools/cuda-ab/inplace_scale.cu @@ -0,0 +1,86 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * inplace_scale.cu - minimal in-place scale repro for the streams2 + * "scale wrote 1 page" class. + * + * One H2D, one in-place scale launch (c[i] *= 2 over 128 KB = 32 pages), + * one D2H, per-page diff. Run it repeatedly: if the launch writes only + * part of the buffer the per-page report names how many 4 KB blocks + * changed and the host log carries the launch's retire wait. + * + * Build: nvcc -O2 --cudart=static -o inplace_scale inplace_scale.cu + */ + +#include <cstdio> +#include <cstdlib> +#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) + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +int main(void) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *h, *d; + int i, bad = 0, first = -1, last = -1; + int nb = (n + 1023) / 1024, blk_bad = 0; + + setvbuf(stdout, NULL, _IONBF, 0); + h = (float *)malloc(bytes); + if (!h) + return 1; + for (i = 0; i < n; i++) + h[i] = (float)(i % 1000) * 0.5f; + + CHECK(cudaMalloc(&d, bytes)); + CHECK(cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice)); + printf("PTR d=%p\n", (void *)d); + + scale<<<(n + 255) / 256, 256>>>(d, 2.0f, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost)); + for (i = 0; i < n; i++) { + float exp = (float)(i % 1000) * 0.5f * 2.0f; + + if (h[i] != exp) { + bad++; + if (first < 0) + first = i; + last = i; + } + } + for (i = 0; i < nb; i++) { + int j, hit = 0; + + for (j = i * 1024; j < (i + 1) * 1024 && j < n; j++) + if (h[j] != (float)(j % 1000) * 0.5f * 2.0f) { + hit = 1; + break; + } + if (hit) + blk_bad++; + } + printf("SCALE n=%d bad=%d first=%d last=%d blocks=%d/%d\n", + n, bad, first, last, blk_bad, nb); + printf("STATUS %s\n", bad ? "FAIL" : "OK"); + return bad ? 1 : 0; +} diff --git a/tools/cuda-ab/ip_probe.cu b/tools/cuda-ab/ip_probe.cu new file mode 100644 index 0000000..366fbc8 --- /dev/null +++ b/tools/cuda-ab/ip_probe.cu @@ -0,0 +1,128 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe.cu - how many CTAs actually executed for an in-place kernel? + * + * scale_mark marks progress[blockIdx.x]=0xCAFEBABE after the c store; D2H of + * progress counts executed CTAs exactly. c (the in-place buffer) is read back + * and verified. seed (argv[1]) makes the run content-unique. + * + * Build: nvcc -O2 --cudart=static -o ip_probe ip_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#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) + +__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]; +} + +__global__ void scale_mark(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xCAFEBABEu; +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[128]; + int i; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g grid=%d\n", prop.name, (double)k, grid); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned))); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + /* in-place scale with CTA markers */ + scale_mark<<<grid, 256>>>(dc, 2.0f, n, d_prog); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int marked = 0, first_missing = -1; + + for (i = 0; i < grid; i++) { + if (h_prog[i] == 0xCAFEBABEu) + marked++; + else if (first_missing < 0) + first_missing = i; + } + printf("CTA-MARK marked=%d/%d first_missing=%d\n", marked, grid, + first_missing); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * 2.0f) { + b++; + if (first < 0) + first = i; + } + printf("SCALE readback bad=%d/%d first=%d\n", b, n, first); + if (first >= 0) + printf("SCALE firstbad got=%g want=%g\n", hc[first], + (ha[first] + hb[first]) * 2.0f); + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/ip_probe2.cu b/tools/cuda-ab/ip_probe2.cu new file mode 100644 index 0000000..a5fa2ec --- /dev/null +++ b/tools/cuda-ab/ip_probe2.cu @@ -0,0 +1,201 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe2.cu - why do in-place c stores vanish while prog stores land? + * + * After a vadd that fills c, three variants run on the SAME c: + * 1. scale_if : if (i<n) c[i] *= s; (control, like ip_probe) + * 2. scale_u : c[i] *= s; (unconditional, grid exact) + * 3. fill_c : c[i] = 7.0f; (store only, no load) + * 4. add_c : c[i] = c[i] + 1.0f; (load + store) + * Each stage reads c back and compares against the expected value at that + * stage. A separate prog buffer marks CTA execution. + * + * Build: nvcc -O2 --cudart=static -o ip_probe2 ip_probe2.cu + */ + +#include <cstdio> +#include <cstdlib> +#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) + +__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]; +} + +__global__ void scale_if(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xCAFEBABEu; +} + +__global__ void scale_u(float *c, float s, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xDEADBEEFu; +} + +__global__ void fill_c(float *c, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = 7.0f; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xFEEDFACu; +} + +__global__ void add_c(float *c, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = c[i] + 1.0f; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0x0DDF00Du; +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[128]; + int i, stage; + float mult; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g\n", prop.name, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + mult = 2.0f; + for (stage = 0; stage < 4; stage++) { + unsigned mark = 0; + + CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned))); + switch (stage) { + case 0: + scale_if<<<grid, 256>>>(dc, 2.0f, n, d_prog); + mult = 2.0f; + mark = 0xCAFEBABEu; + break; + case 1: + scale_u<<<grid, 256>>>(dc, 2.0f, d_prog); + mult = 4.0f; + mark = 0xDEADBEEFu; + break; + case 2: + fill_c<<<grid, 256>>>(dc, d_prog); + mark = 0xFEEDFACu; + break; + default: + add_c<<<grid, 256>>>(dc, d_prog); + mult = -1.0f; + mark = 0x0DDF00Du; + break; + } + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int marked = 0; + + for (i = 0; i < grid; i++) + if (h_prog[i] == mark) + marked++; + printf("S%d CTA-MARK marked=%d/%d\n", stage + 1, + marked, grid); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + if (stage < 2) { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * mult) { + b++; + if (first < 0) + first = i; + } + printf("S%d readback bad=%d/%d first=%d\n", + stage + 1, b, n, first); + } else if (stage == 2) { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != 7.0f) { + b++; + if (first < 0) + first = i; + } + printf("S3 readback bad=%d/%d first=%d (fill 7.0)\n", + b, n, first); + } else { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != 8.0f) { + b++; + if (first < 0) + first = i; + } + printf("S4 readback bad=%d/%d first=%d (add 1.0 to fill)\n", + b, n, first); + } + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/ip_probe3.cu b/tools/cuda-ab/ip_probe3.cu new file mode 100644 index 0000000..a50df0e --- /dev/null +++ b/tools/cuda-ab/ip_probe3.cu @@ -0,0 +1,142 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe3.cu - what does the SM actually READ for the scale kernel? + * + * scale_report writes, per CTA: prog[2b] = bits of the float param s, + * prog[2b+1] = bits of c[b*256] (the load). + * The D2H of prog then shows exactly what the SM saw, independent of any + * store-visibility path for c. + * + * Build: nvcc -O2 --cudart=static -o ip_probe3 ip_probe3.cu + */ + +#include <cstdio> +#include <cstdlib> +#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) + +__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]; +} + +__global__ void scale_report(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) { + prog[2 * blockIdx.x] = __float_as_uint(s); + prog[2 * blockIdx.x + 1] = __float_as_uint(c[blockIdx.x * 256]); + } +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[256]; + int i; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g\n", prop.name, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 256 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_prog, 0xEE, 256 * sizeof(unsigned))); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + scale_report<<<grid, 256>>>(dc, 2.0f, n, d_prog); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 256 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int s_ok = 0, s_1 = 0, c_ok = 0, n_seen = 0; + unsigned first_s = 0, first_c = 0; + + for (i = 0; i < grid; i++) { + unsigned sb = h_prog[2 * i]; + unsigned cb = h_prog[2 * i + 1]; + + if (i == 0) { + first_s = sb; + first_c = cb; + } + if (sb == 0x40000000u) + s_ok++; + if (sb == 0x3F800000u) + s_1++; + { + float want = ha[i * 256] + hb[i * 256]; + + if (cb == *(unsigned *)&want) + c_ok++; + } + n_seen++; + } + printf("REPORT ctas=%d s_eq_2.0=%d s_eq_1.0=%d c_load_ok=%d first_s=0x%08x first_c=0x%08x\n", + n_seen, s_ok, s_1, c_ok, first_s, first_c); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * 2.0f) { + b++; + if (first < 0) + first = i; + } + printf("SCALE readback bad=%d/%d first=%d\n", b, n, first); + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/launch_count.cu b/tools/cuda-ab/launch_count.cu new file mode 100644 index 0000000..02db1d6 --- /dev/null +++ b/tools/cuda-ab/launch_count.cu @@ -0,0 +1,25 @@ +// SPDX-License-Identifier: BSD-3-Clause +// launch_count.cu - single-launch execution-count repro (r342). +// +// One launch, one atomic increment. The engine-written page witness showed the +// increment applied TWICE (1 -> 2; the 100-increment variant -> 200): every +// forwarded launch executes its engine work twice (one submission, one ring +// publish, completion OK). This is the root cause of the multi-pass content +// class (victory_reduce exactly 2x, scale/saxpy/smiley/multikernel wrong). +// Run in the guest; expected LAUNCH_COUNT=1. +// for a single launch, the launch path executes work more than once. +#include <cstdio> +#include <cuda_runtime.h> +__global__ void inc(unsigned long long *c) { if (threadIdx.x == 0 && blockIdx.x == 0) atomicAdd(c, 1ULL); } +int main(void) { + unsigned long long *d; unsigned long long h = 0; + if (cudaMalloc(&d, 8) != cudaSuccess) { printf("CUDA_ERROR malloc\nSTATUS FAIL\n"); return 1; } + if (cudaMemset(d, 0, 8) != cudaSuccess) { printf("CUDA_ERROR memset\nSTATUS FAIL\n"); return 1; } + inc<<<1, 32>>>(d); + if (cudaGetLastError() != cudaSuccess) { printf("CUDA_ERROR launch\nSTATUS FAIL\n"); return 1; } + if (cudaDeviceSynchronize() != cudaSuccess) { printf("CUDA_ERROR sync\nSTATUS FAIL\n"); return 1; } + if (cudaMemcpy(&h, d, 8, cudaMemcpyDeviceToHost) != cudaSuccess) { printf("CUDA_ERROR d2h\nSTATUS FAIL\n"); return 1; } + printf("LAUNCH_COUNT=%llu (expect 1)\n", h); + printf("STATUS %s\n", h == 1 ? "OK" : "FAIL"); + return h == 1 ? 0 : 1; +} diff --git a/tools/cuda-ab/memset2d_probe.cu b/tools/cuda-ab/memset2d_probe.cu new file mode 100644 index 0000000..acf43d4 --- /dev/null +++ b/tools/cuda-ab/memset2d_probe.cu @@ -0,0 +1,157 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memset2d_probe.cu - per-phase diagnostics for victory_memset2d. + * + * victory_memset2d reports one aggregate bad count (~12.5M) that mixes all + * three phases. This probe runs the SAME three phases and reports each + * phase's bad count, the first mismatching byte (offset, got, want), and a + * per-phase 1MB-chunk histogram so the failing phase and the failing region + * are attributable in one run. + * + * Phases: + * P1 memset : cudaMemset(d_img, 0xA5, pitch*h); D2H 2D readback + * P2 roundtrip: H2D 2D pattern; D2H 2D readback + * P3 memset2d : cudaMemset2D quadrant 0x3C; D2H 2D readback + * + * Output lines: + * DEVICE ... + * PHASE p=1 bad=<n> first_off=<i> got=<g> want=<w> + * PHIST p=1 mb=<m> ok=<n> bad=<n> + * MEMSET2D2 bad_total=<n> + * STATUS OK|FAIL + */ + +#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 IMG_W 4096 +#define IMG_H 1024 +#define FILL_BYTE 0xA5 +#define QUAD_BYTE 0x3C + +static int phase_check(const char *tag, const unsigned char *back, + const unsigned char *want, size_t n, int p) +{ + size_t i, first = (size_t)-1; + int bad = 0, got = 0, wantv = 0; + size_t mb; + + for (i = 0; i < n; i++) { + if (back[i] != want[i]) { + if (first == (size_t)-1) { + first = i; + got = back[i]; + wantv = want[i]; + } + bad++; + } + } + printf("PHASE p=%d bad=%d first_off=%zu got=0x%02x want=0x%02x\n", + p, bad, first, got, wantv); + if (bad) { + size_t nb = (n + (1u << 20) - 1) >> 20; + + for (mb = 0; mb < nb && mb < 8; mb++) { + size_t lo = mb << 20, hi = lo + (1u << 20); + size_t j; + int bok = 0, bbad = 0; + + if (hi > n) + hi = n; + for (j = lo; j < hi; j++) { + if (back[j] == want[j]) + bok++; + else + bbad++; + } + printf("PHIST p=%d mb=%zu ok=%d bad=%d\n", + p, mb, bok, bbad); + } + } + (void)tag; + return bad; +} + +int main(void) +{ + cudaDeviceProp prop; + const size_t w = IMG_W, h = IMG_H; + unsigned char *h_img, *h_back, *h_want; + unsigned char *d_img; + size_t pitch = 0; + int bad = 0, row, col; + + 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_img = (unsigned char *)malloc(w * h); + h_back = (unsigned char *)malloc(w * h); + h_want = (unsigned char *)malloc(w * h); + if (!h_img || !h_back || !h_want) + return 1; + + CHECK(cudaMallocPitch(&d_img, &pitch, w, h)); + printf("PITCH pitch=%zu\n", pitch); + printf("DPTRS h_img=%p h_back=%p h_want=%p d_img=%p\n", (void *)h_img, + (void *)h_back, (void *)h_want, (void *)d_img); + + /* P1: full cudaMemset over the pitched allocation, 2D readback. */ + CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + memset(h_want, FILL_BYTE, w * h); + bad += phase_check("P1", h_back, h_want, w * h, 1); + + /* P2: pitched H2D pattern -> D2H round trip. */ + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) + h_img[row * w + col] = (unsigned char)((row + col) & 0xFF); + CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h, + cudaMemcpyHostToDevice)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + bad += phase_check("P2", h_back, h_img, w * h, 2); + + /* P3: cudaMemset2D of the upper-left quadrant only. */ + CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2)); + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) { + unsigned char expect = h_img[row * w + col]; + + if (row < (int)(h / 2) && col < (int)(w / 2)) + expect = QUAD_BYTE; + h_want[row * w + col] = expect; + } + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + bad += phase_check("P3", h_back, h_want, w * h, 3); + + printf("MEMSET2D2 bad_total=%d\n", bad); + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/memset_probe.cu b/tools/cuda-ab/memset_probe.cu new file mode 100644 index 0000000..a6a12e0 --- /dev/null +++ b/tools/cuda-ab/memset_probe.cu @@ -0,0 +1,87 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memset_probe.cu - isolate the memcpy_roundtrip memset-phase failure. + * + * cudaMemset(8 MB, 0x5A) -> D2H -> byte verify, with a run-length report: + * the first bad offset, and a coverage map at 4 KB granularity (correct / + * wrong / untouched-malloc-prefill) so "fill under-delivered" is + * distinguishable from "readback delivered wrong pages". + * + * Build: nvcc -O2 --cudart=static -o memset_probe memset_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#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 PAGE 4096 +#define MemsetByte 0x5A + +int main(void) +{ + const size_t bytes = 8u << 20; + unsigned char *h; + unsigned char *d; + size_t i; + int bad = 0; + size_t first_bad = (size_t)-1; + size_t n_pages, p; + + setvbuf(stdout, NULL, _IONBF, 0); + h = (unsigned char *)malloc(bytes); + if (!h) + return 1; + memset(h, 0xEE, bytes); + + CHECK(cudaMalloc(&d, bytes)); + printf("DPTR d=%p bytes=%zu\n", (void *)d, bytes); + + CHECK(cudaMemset(d, MemsetByte, bytes)); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < bytes; i++) { + if (h[i] != (unsigned char)MemsetByte) { + if (first_bad == (size_t)-1) + first_bad = i; + bad++; + } + } + printf("FIRST_BAD off=%zu (0x%zx) bad=%d total=%zu\n", + first_bad, first_bad, bad, bytes); + n_pages = bytes / PAGE; + for (p = 0; p < n_pages; p++) { + int good = 1, prefill = 1; + + for (i = 0; i < PAGE; i++) { + if (h[p * PAGE + i] != (unsigned char)MemsetByte) + good = 0; + if (h[p * PAGE + i] != 0xEE) + prefill = 0; + if (!good && !prefill) + break; + } + if (!good) { + printf("PAGE %zu %s b0=%02x b1=%02x b2=%02x b3=%02x b4095=%02x\n", + p, prefill ? "PREFILL" : "OTHER", + h[p * PAGE], h[p * PAGE + 1], + h[p * PAGE + 2], h[p * PAGE + 3], + h[p * PAGE + 4095]); + if (p > 24) + break; + } + } + printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL"); + return 0; +} diff --git a/tools/cuda-ab/push_battery.sh b/tools/cuda-ab/push_battery.sh new file mode 100755 index 0000000..7bd3df9 --- /dev/null +++ b/tools/cuda-ab/push_battery.sh @@ -0,0 +1,22 @@ +#!/bin/bash +# push_battery.sh - push every CUDA battery binary from tools/cuda-ab/run/ +# to the running guest via scp over the SSH hostfwd (port 2222). +# Usage: push_battery.sh [guest_dir] (default /tmp) +set -u +HERE="$(cd "$(dirname "$0")" && pwd)" +RUN="$HERE/run" +GUEST_DIR="${1:-/tmp}" +CLIENTS="victory_mail smiley vecadd_loop streams2 memcpy_roundtrip multikernel victory_saxpy victory_reduce victory_matmul victory_strided victory_memset2d victory_atomics victory_streams_events victory_math vector_add inplace_scale streams2_probe streams_events_probe memset2d_probe tgemm_probe transformer transformer_stages ip_probe ip_probe2 ip_probe3 scale_probe rt_probe launch_count egl_frame frame_producer" +export SSHPASS="${VGPU_GUEST_PASS:-ubuntu}" +FILES="" +for c in $CLIENTS; do [ -x "$RUN/$c" ] && FILES="$FILES $RUN/$c"; done +[ -n "$FILES" ] || { echo "no binaries in $RUN"; exit 1; } +# shellcheck disable=SC2086 +sshpass -e scp -P "${VGPU_GUEST_PORT:-2222}" \ + -o StrictHostKeyChecking=no -o UserKnownHostsFile=/dev/null \ + -o PubkeyAuthentication=no -o PreferredAuthentications=password \ + -o ConnectTimeout=15 $FILES "ubuntu@127.0.0.1:$GUEST_DIR/" && \ +sshpass -e ssh -p "${VGPU_GUEST_PORT:-2222}" \ + -o StrictHostKeyChecking=no -o UserKnownHostsFile=/dev/null \ + -o PubkeyAuthentication=no -o PreferredAuthentications=password \ + "ubuntu@127.0.0.1" "chmod +x $GUEST_DIR/* 2>/dev/null; ls $GUEST_DIR | wc -l; echo BATTERY_PUSHED" diff --git a/tools/cuda-ab/rt_probe.cu b/tools/cuda-ab/rt_probe.cu new file mode 100644 index 0000000..025301d --- /dev/null +++ b/tools/cuda-ab/rt_probe.cu @@ -0,0 +1,165 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * rt_probe.cu - instrumented memcpy_roundtrip variant isolating the + * memset-phase D2H readback failure. + * + * Phases (argv[1]): + * "full" : H2D -> kernel -> D2H -> memset -> D2H (the original order) + * "sync" : same but cudaDeviceSynchronize() before the second D2H + * "memset" : memset -> D2H only (memset_probe shape, 3-buffer alloc) + * "rev" : memset -> D2H first, then H2D -> kernel -> D2H + * + * Forensics on the memset readback: first bad offset, byte histogram + * (0x5A / 0x00 / other), per-1MB chunk histogram, and the first 32 bytes. + * + * Build: nvcc -O2 --cudart=static -o rt_probe rt_probe.cu + */ + +#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; +} + +static void memset_forensics(const unsigned char *h, size_t bytes, + const char *tag) +{ + size_t i, first_bad = (size_t)-1; + int bad = 0, n5a = 0, n00 = 0, nother = 0; + size_t chunk; + + for (i = 0; i < bytes; i++) { + if (h[i] == 0x5A) + n5a++; + else if (h[i] == 0x00) + n00++; + else + nother++; + if (h[i] != 0x5A) { + if (first_bad == (size_t)-1) + first_bad = i; + bad++; + } + } + printf("MEMSET-READ %s bad=%d first_bad=%zu n5a=%d n00=%d nother=%d\n", + tag, bad, first_bad, n5a, n00, nother); + printf("MEMSET-HEAD %s", tag); + for (i = 0; i < 32; i++) + printf(" %02x", h[i]); + printf("\n"); + if (first_bad != (size_t)-1) { + printf("MEMSET-ATBAD %s off=%zu:", tag, first_bad); + for (i = first_bad; i < first_bad + 16 && i < bytes; i++) + printf(" %02x", h[i]); + printf("\n"); + } + for (chunk = 0; chunk < bytes / (1u << 20); chunk++) { + size_t lo = chunk << 20, hi = lo + (1u << 20); + int c5a = 0, c00 = 0, cother = 0; + + for (i = lo; i < hi; i++) { + if (h[i] == 0x5A) + c5a++; + else if (h[i] == 0x00) + c00++; + else + cother++; + } + printf("MEMSET-MB %s mb=%zu 5a=%d 00=%d other=%d\n", + tag, chunk, c5a, c00, cother); + } +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "full"; + cudaDeviceProp prop; + const size_t bytes = 8u << 20; + 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 mode=%s\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem, mode); + + 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)); + printf("DPTRS d_in=%p d_out=%p d_m=%p\n", (void *)d_in, + (void *)d_out, (void *)d_m); + + if (!strcmp(mode, "memset") || !strcmp(mode, "rev")) { + /* memset phase first */ + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + if (!strcmp(mode, "sync")) + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + memset_forensics(h_m, bytes, "early"); + } + + if (!strcmp(mode, "full") || !strcmp(mode, "sync") || + !strcmp(mode, "rev")) { + 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]; + } + printf("BUMP bad=%d csum=%.6f\n", bad, csum); + } + + if (!strcmp(mode, "full") || !strcmp(mode, "sync")) { + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + if (!strcmp(mode, "sync")) + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + memset_forensics(h_m, bytes, "late"); + } + + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/run/egl_frame b/tools/cuda-ab/run/egl_frame Binary files differnew file mode 100755 index 0000000..130f032 --- /dev/null +++ b/tools/cuda-ab/run/egl_frame diff --git a/tools/cuda-ab/run/evhold b/tools/cuda-ab/run/evhold Binary files differnew file mode 100755 index 0000000..de4448c --- /dev/null +++ b/tools/cuda-ab/run/evhold diff --git a/tools/cuda-ab/run/frame_producer b/tools/cuda-ab/run/frame_producer Binary files differnew file mode 100755 index 0000000..a202ac7 --- /dev/null +++ b/tools/cuda-ab/run/frame_producer diff --git a/tools/cuda-ab/run/inplace_scale b/tools/cuda-ab/run/inplace_scale Binary files differnew file mode 100755 index 0000000..bb47f42 --- /dev/null +++ b/tools/cuda-ab/run/inplace_scale diff --git a/tools/cuda-ab/run/ip4 b/tools/cuda-ab/run/ip4 Binary files differnew file mode 100755 index 0000000..54937e6 --- /dev/null +++ b/tools/cuda-ab/run/ip4 diff --git a/tools/cuda-ab/run/ip5 b/tools/cuda-ab/run/ip5 Binary files differnew file mode 100755 index 0000000..c4b4a39 --- /dev/null +++ b/tools/cuda-ab/run/ip5 diff --git a/tools/cuda-ab/run/ip6 b/tools/cuda-ab/run/ip6 Binary files differnew file mode 100755 index 0000000..56815c2 --- /dev/null +++ b/tools/cuda-ab/run/ip6 diff --git a/tools/cuda-ab/run/ip_probe b/tools/cuda-ab/run/ip_probe Binary files differnew file mode 100755 index 0000000..08fb86e --- /dev/null +++ b/tools/cuda-ab/run/ip_probe diff --git a/tools/cuda-ab/run/ip_probe2 b/tools/cuda-ab/run/ip_probe2 Binary files differnew file mode 100755 index 0000000..dfe6ab9 --- /dev/null +++ b/tools/cuda-ab/run/ip_probe2 diff --git a/tools/cuda-ab/run/ip_probe3 b/tools/cuda-ab/run/ip_probe3 Binary files differnew file mode 100755 index 0000000..8a34c05 --- /dev/null +++ b/tools/cuda-ab/run/ip_probe3 diff --git a/tools/cuda-ab/run/launch_count b/tools/cuda-ab/run/launch_count Binary files differnew file mode 100755 index 0000000..0c67f45 --- /dev/null +++ b/tools/cuda-ab/run/launch_count diff --git a/tools/cuda-ab/run/memset2d_probe b/tools/cuda-ab/run/memset2d_probe Binary files differnew file mode 100755 index 0000000..f5bd071 --- /dev/null +++ b/tools/cuda-ab/run/memset2d_probe diff --git a/tools/cuda-ab/run/memset_probe b/tools/cuda-ab/run/memset_probe Binary files differnew file mode 100755 index 0000000..bd9df5e --- /dev/null +++ b/tools/cuda-ab/run/memset_probe diff --git a/tools/cuda-ab/run/rt_probe b/tools/cuda-ab/run/rt_probe Binary files differnew file mode 100755 index 0000000..48d866f --- /dev/null +++ b/tools/cuda-ab/run/rt_probe diff --git a/tools/cuda-ab/run/scale_probe b/tools/cuda-ab/run/scale_probe Binary files differnew file mode 100755 index 0000000..ca3fc20 --- /dev/null +++ b/tools/cuda-ab/run/scale_probe diff --git a/tools/cuda-ab/run/streams2_probe b/tools/cuda-ab/run/streams2_probe Binary files differnew file mode 100755 index 0000000..206d6d6 --- /dev/null +++ b/tools/cuda-ab/run/streams2_probe diff --git a/tools/cuda-ab/run/streams2d b/tools/cuda-ab/run/streams2d Binary files differnew file mode 100755 index 0000000..90b6df8 --- /dev/null +++ b/tools/cuda-ab/run/streams2d diff --git a/tools/cuda-ab/run/streams_events_probe b/tools/cuda-ab/run/streams_events_probe Binary files differnew file mode 100755 index 0000000..00f3539 --- /dev/null +++ b/tools/cuda-ab/run/streams_events_probe diff --git a/tools/cuda-ab/run/tgemm_probe b/tools/cuda-ab/run/tgemm_probe Binary files differnew file mode 100755 index 0000000..d9c9aa3 --- /dev/null +++ b/tools/cuda-ab/run/tgemm_probe diff --git a/tools/cuda-ab/run/transformer b/tools/cuda-ab/run/transformer Binary files differnew file mode 100755 index 0000000..bcde39e --- /dev/null +++ b/tools/cuda-ab/run/transformer diff --git a/tools/cuda-ab/run/transformer_async b/tools/cuda-ab/run/transformer_async Binary files differnew file mode 100755 index 0000000..2911b5c --- /dev/null +++ b/tools/cuda-ab/run/transformer_async diff --git a/tools/cuda-ab/run/transformer_stages b/tools/cuda-ab/run/transformer_stages Binary files differnew file mode 100755 index 0000000..1461a3e --- /dev/null +++ b/tools/cuda-ab/run/transformer_stages diff --git a/tools/cuda-ab/run/wedge_ab b/tools/cuda-ab/run/wedge_ab Binary files differnew file mode 100755 index 0000000..23a90aa --- /dev/null +++ b/tools/cuda-ab/run/wedge_ab diff --git a/tools/cuda-ab/run/wedge_bisect b/tools/cuda-ab/run/wedge_bisect Binary files differnew file mode 100755 index 0000000..5898fcf --- /dev/null +++ b/tools/cuda-ab/run/wedge_bisect diff --git a/tools/cuda-ab/run/wedge_recover b/tools/cuda-ab/run/wedge_recover Binary files differnew file mode 100755 index 0000000..3d9351c --- /dev/null +++ b/tools/cuda-ab/run/wedge_recover diff --git a/tools/cuda-ab/scale_probe.cu b/tools/cuda-ab/scale_probe.cu new file mode 100644 index 0000000..8a0df59 --- /dev/null +++ b/tools/cuda-ab/scale_probe.cu @@ -0,0 +1,245 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * scale_probe.cu - localize the deterministic "in-place scale writes only the + * first 4 KB page" failure. + * + * Modes (argv[1]): + * V : vadd(a,b->c); sync; D2H+verify; vadd again; D2H+verify + * I : vadd(a,b->c); sync; D2H+verify; scale(c,*2 in place); sync; + * D2H+verify (with page map); scale again; sync; D2H+verify; D2H again + * O : vadd(a,b->c); sync; D2H+verify; scale_out(c->d,*2); sync; + * D2H+verify (with page map) + * + * Every stage also runs a GPU-side verify kernel that counts matching + * elements reading the device buffers (engine's own view) and writes the + * counts to a small buffer, D2H'd afterwards: separates "device pages never + * written" (GPU count low) from "D2H readback stale" (GPU count high, host + * count low). + * + * Build: nvcc -O2 --cudart=static -o scale_probe scale_probe.cu + */ + +#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) + +__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]; +} + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +__global__ void scale_out(const float *c, float *d, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + d[i] = c[i] * s; +} + +/* verify kernel: out[0]=count(c==a+b), out[1]=count(c==2*(a+b)) */ +__global__ void vdigest(const float *a, const float *b, const float *c, int n, + unsigned long long *out) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + if (c[i] == a[i] + b[i]) + atomicAdd(&out[0], 1ull); + if (c[i] == (a[i] + b[i]) * 2.0f) + atomicAdd(&out[1], 1ull); +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "V"; + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc, *hd; + float *da, *db, *dc, *dd; + unsigned long long *d_out, h_out[2]; + int i; + float k = argc > 2 ? atof(argv[2]) : 0.0f; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + { + FILE *f = fopen("/proc/uptime", "r"); + double up = -1.0; + + if (f) { if (fscanf(f, "%lf", &up) != 1) up = -1.0; fclose(f); } + printf("UPTIME %.3f\n", up); + } + printf("DEVICE %s mode=%s seed=%g\n", prop.name, mode, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + hd = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + memset(hd, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&dd, bytes)); + CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long))); + printf("DPTRS da=%p db=%p dc=%p dd=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)dd); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("S1 vadd readback bad=%d/%d\n", b, n); + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dc, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, + 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S1 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", h_out[0], + h_out[1]); + } + + if (!strcmp(mode, "V")) { + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("S2 vadd-again readback bad=%d/%d\n", b, n); + } + printf("STATUS DONE\n"); + return 0; + } + + if (!strcmp(mode, "I")) { + int stage; + + for (stage = 0; stage < 2; stage++) { + float mult = stage == 0 ? 2.0f : 4.0f; + + scale<<<grid, 256>>>(dc, 2.0f, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, page; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * mult) + b++; + printf("S%d inplace-scale readback bad=%d/%d (mult=%g)\n", + 2 + stage, b, n, mult); + if (stage == 0) + for (page = 0; page < 32; page++) { + int pb = 0, k; + + for (k = 0; k < 256; k++) + if (hc[page * 256 + k] != + (ha[page * 256 + k] + + hb[page * 256 + k]) * mult) + pb++; + if (pb) + printf("S%d PAGE %d bad=%d/256 first=%g want=%g\n", + 2 + stage, page, + pb, hc[page * 256], + (ha[page * 256] + + hb[page * 256]) * mult); + } + } + if (stage == 0) { + CHECK(cudaMemset(d_out, 0, + 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dc, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, + 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S2 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", + h_out[0], h_out[1]); + } + } + printf("STATUS DONE\n"); + return 0; + } + + if (!strcmp(mode, "O")) { + scale_out<<<grid, 256>>>(dc, dd, 2.0f, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hd, dd, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, page; + + for (i = 0; i < n; i++) + if (hd[i] != (ha[i] + hb[i]) * 2.0f) + b++; + printf("S2 outofplace-scale readback bad=%d/%d\n", b, n); + for (page = 0; page < 32; page++) { + int pb = 0, k; + + for (k = 0; k < 256; k++) + if (hd[page * 256 + k] != + (ha[page * 256 + k] + + hb[page * 256 + k]) * 2.0f) + pb++; + if (pb) + printf("S2 PAGE %d bad=%d/256\n", page, + pb); + } + } + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dd, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S2 GPU-DIGEST d_eq_ab=%llu d_eq_2ab=%llu\n", h_out[0], + h_out[1]); + printf("STATUS DONE\n"); + return 0; + } + return 1; +} diff --git a/tools/cuda-ab/soak.sh b/tools/cuda-ab/soak.sh index f870fda..4f1430f 100755 --- a/tools/cuda-ab/soak.sh +++ b/tools/cuda-ab/soak.sh @@ -135,6 +135,19 @@ run_one() { r="$(gssh "cat $out 2>/dev/null")" || true [ -n "$r" ] && text="$r" printf '%s' "$text" | grep -q 'EXIT=' && break + # The guest wrapper's EXIT= marker can be lost when the guest's + # process reaping stalls (observed: the client reached STATUS OK + # but its zombie was never reaped and `timeout` never wrote the + # marker). A terminal STATUS line is proof of completion by + # itself - accept it after a short grace so the EXIT= marker + # still wins when it does arrive. + if printf '%s' "$text" | grep -qE 'STATUS (OK|FAIL)'; then + sleep 10 + r="$(gssh "cat $out 2>/dev/null")" || true + [ -n "$r" ] && text="$r" + printf '%s' "$text" | grep -q 'EXIT=' && break + break + fi tnow=$(date +%s) if [ $((tnow - t0)) -gt $((TMO + 120)) ]; then if client_alive "$bin"; then @@ -179,7 +192,7 @@ log "SOAK START mode=$MODE seconds=$SECS guest_dir=$GUEST_DIR timeout=${TMO}s lo # cannot silently shrink the battery. found="" for i in 1 2 3; do - found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done")" || found="" + found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done; true")" || found="" [ -n "$found" ] && break sleep 3 done diff --git a/tools/cuda-ab/streams2_probe.cu b/tools/cuda-ab/streams2_probe.cu new file mode 100644 index 0000000..5bebab4 --- /dev/null +++ b/tools/cuda-ab/streams2_probe.cu @@ -0,0 +1,299 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2_probe.cu - instrumented streams2 to localize the deterministic + * failure: per-array mismatch counts, first-bad indices with values, and a + * GPU-side verify-kernel digest of the device buffers (to separate "device + * pages never written" from "D2H readback wrong"). + * + * Modes (argv[1]): + * A : exactly streams2's shape (control) + * B : like A but the two D2H copy-backs use synchronous cudaMemcpy + * C : like A but no events cross the streams; both streams are joined + * with cudaStreamSynchronize before any copy-back + * D : single stream (everything on s1) + * + * Build: nvcc -O2 --cudart=static -o streams2_probe streams2_probe.cu + */ + +#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) + +__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]; +} + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +/* GPU-side digest: counts of matches for c1==a1+b1 and c2==(a2+b2)*2 */ +__global__ void digest(const float *a1, const float *b1, const float *c1, + const float *a2, const float *b2, const float *c2, + int n, unsigned long long *out) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + if (c1[i] == a1[i] + b1[i]) + atomicAdd(&out[0], 1ull); + if (c2[i] == (a2[i] + b2[i]) * 2.0f) + atomicAdd(&out[1], 1ull); +} + +static int report(const char *tag, const float *h_c, const float *h_a, + const float *h_b, const float *h_a2, const float *h_b2, + int n, float scale2) +{ + int bad = 0, i, shown = 0; + + for (i = 0; i < n; i++) { + float want = h_a[i] + h_b[i]; + + if (scale2 != 1.0f) + want = (h_a2[i] + h_b2[i]) * scale2; + if (h_c[i] != want) { + bad++; + if (shown < 4) { + printf("FIRSTBAD %s i=%d got=%g want=%g\n", + tag, i, h_c[i], want); + shown++; + } + } + } + printf("BAD %s n=%d bad=%d\n", tag, n, bad); + return bad; +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "A"; + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + unsigned long long *d_out, h_out[2]; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + int i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s mode=%s\n", prop.name, mode); + + h_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + memset(h_c1, 0xCD, bytes); + memset(h_c2, 0xCD, bytes); + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long))); + printf("DPTRS da1=%p db1=%p dc1=%p da2=%p db2=%p dc2=%p\n", + (void *)d_a1, (void *)d_b1, (void *)d_c1, (void *)d_a2, + (void *)d_b2, (void *)d_c2); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + if (!strcmp(mode, "D")) { + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, + cudaMemcpyHostToDevice, s1)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s1>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaStreamSynchronize(s1)); + scale<<<grid, 256, 0, s1>>>(d_c2, s2, n); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost)); + goto verify; + } + + if (!strcmp(mode, "S")) { + /* staged verification: check device content after every step */ + float *hb1 = (float *)malloc(bytes); + float *hb2 = (float *)malloc(bytes); + float *hb3 = (float *)malloc(bytes); + + CHECK(cudaMemcpy(d_a1, h_a1, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b1, h_b1, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_a2, h_a2, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b2, h_b2, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(hb1, d_a1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a1[i]) b++; + printf("STAGE-IN da1 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_b1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b1[i]) b++; + printf("STAGE-IN db1 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_a2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a2[i]) b++; + printf("STAGE-IN da2 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_b2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b2[i]) b++; + printf("STAGE-IN db2 bad=%d\n", b); } + + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb1[i] != h_a1[i] + h_b1[i]) { b++; + if (first < 0) { first = i; + printf("S1-FIRSTBAD i=%d got=%g want=%g a2b2=%g\n", + i, hb1[i], h_a1[i]+h_b1[i], h_a2[i]+h_b2[i]); } } + printf("STAGE vadd1->c1 readback bad=%d first=%d\n", b, first); } + + CHECK(cudaDeviceSynchronize()); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb2, d_c2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb2[i] != h_a2[i] + h_b2[i]) { b++; if (first < 0) first = i; } + printf("STAGE vadd2->c2 readback bad=%d first=%d\n", b, first); } + + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb3, d_c2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb3[i] != (h_a2[i] + h_b2[i]) * 2.0f) { b++; if (first < 0) first = i; } + printf("STAGE scale->c2 readback bad=%d first=%d\n", b, first); } + /* second readback of c1 AFTER everything: does it change? */ + CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb1[i] != h_a1[i] + h_b1[i]) { b++; if (first < 0) first = i; } + printf("STAGE c1-late readback bad=%d first=%d\n", b, first); } + printf("STATUS OK\n"); + return 0; + } + + if (!strcmp(mode, "C")) { + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, + cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, + cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, + cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, + cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + goto verify; + } + + /* A and B: original shape */ + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + if (!strcmp(mode, "B")) { + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost)); + } else { + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, + cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, + cudaMemcpyDeviceToHost, s2s)); + } + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + +verify: + { + int bad1, bad2; + + bad1 = report("C1", h_c1, h_a1, h_b1, h_a2, h_b2, n, 1.0f); + bad2 = report("C2", h_c2, h_a2, h_b2, h_a2, h_b2, n, s2); + printf("STREAMS n=%d bad=%d\n", n, bad1 + bad2); + printf("TIMING ms=%.3f\n", ms); + + /* GPU-side digest of the device buffers */ + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + digest<<<grid, 256>>>(d_a1, d_b1, d_c1, d_a2, d_b2, d_c2, n, + d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("GPU-DIGEST c1_ok=%llu c2_ok=%llu\n", h_out[0], + h_out[1]); + printf("STATUS %s\n", (bad1 | bad2) ? "FAIL" : "OK"); + } + return 0; +} diff --git a/tools/cuda-ab/streams2d.cu b/tools/cuda-ab/streams2d.cu new file mode 100644 index 0000000..04a0866 --- /dev/null +++ b/tools/cuda-ab/streams2d.cu @@ -0,0 +1,197 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2d.cu - streams2 with per-array / per-range diagnostic reporting. + * + * Same workload as streams2.cu (two streams, cross-stream event ordering) + * but instead of one "bad" count it reports, per array and per 4 KB block, + * how many elements mismatched, and prints the device pointers so a host + * log can be correlated with what the engine read/wrote. + * + * Output lines: + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * PTR <name>=0x<hex> (d_a1 d_b1 d_c1 d_a2 d_b2 d_c2) + * DIFF c1 total=<n> blocks=<first block index with mismatches> + * DIFF c2 total=<n> blocks=<...> + * STREAMS n=<count> bad=<mismatches> + * TIMING ms=<elapsed> + * STATUS OK|FAIL + * + * Build: nvcc -O2 --cudart=static -o streams2d streams2d.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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) + +__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]; +} + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +static void report_diff(const char *name, const float *got, + const float *exp, int n) +{ + int i, bad = 0, first = -1, last = -1; + int blk_hits = 0, blk_first = -1; + int nb = (n + 1023) / 1024; + int b; + + for (i = 0; i < n; i++) { + if (got[i] != exp[i]) { + bad++; + if (first < 0) + first = i; + last = i; + } + } + for (b = 0; b < nb; b++) { + int j, hit = 0; + + for (j = b * 1024; j < (b + 1) * 1024 && j < n; j++) + if (got[j] != exp[j]) { + hit = 1; + break; + } + if (hit) { + blk_hits++; + if (blk_first < 0) + blk_first = b; + } + } + printf("DIFF %s total=%d first=%d last=%d blocks=%d/%d blk_first=%d\n", + name, bad, first, last, blk_hits, nb, blk_first); + if (first >= 0) + printf("DIFF %s sample got[%d]=%a exp[%d]=%a got[%d]=%a exp[%d]=%a\n", + name, first, got[first], first, exp[first], last, + got[last], last, exp[last]); +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + 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_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + if (!h_a1 || !h_b1 || !h_c1 || !h_a2 || !h_b2 || !h_c2) + return 1; + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + printf("PTR d_a1=%p d_b1=%p d_c1=%p d_a2=%p d_b2=%p d_c2=%p\n", + (void *)d_a1, (void *)d_b1, (void *)d_c1, + (void *)d_a2, (void *)d_b2, (void *)d_c2); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + report_diff("c1", h_c1, h_a1, n); + { + float *exp2 = (float *)malloc(bytes); + + for (i = 0; i < n; i++) + exp2[i] = (h_a2[i] + h_b2[i]) * s2; + report_diff("c2", h_c2, exp2, n); + for (i = 0; i < n; i++) { + if (h_c1[i] != h_a1[i] + h_b1[i]) + bad++; + if (h_c2[i] != exp2[i]) + bad++; + } + free(exp2); + } + + printf("STREAMS n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/streams_events_probe.cu b/tools/cuda-ab/streams_events_probe.cu new file mode 100644 index 0000000..94aef69 --- /dev/null +++ b/tools/cuda-ab/streams_events_probe.cu @@ -0,0 +1,170 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams_events_probe.cu - instrumented victory_streams_events. + * + * Same 3-stream event chain (c0 = a0+b0; c1 = (a1+b1)+c0; c2 = (a2+b2)+c1) + * but with per-buffer diagnostics: device pointers, per-buffer bad counts, + * and the first 8 expected/actual values of every output so a stale-input + * vs stale-readback failure is distinguishable in one run. + * + * Output lines: + * DEVICE ... + * DPTRS da0..dc2 (device), ha0..hc2 (host) + * BUF s=<n> bad=<n> + * FIRST s=<n> exp=<v0..v7> got=<v0..v7> + * STREAMS3 n=<count> bad=<mismatches> + * STATUS OK|FAIL + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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 NSTREAMS 3 + +__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]; +} + +__global__ void vadd_into(float *c, const float *s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] + s[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 16; + size_t bytes = (size_t)n * sizeof(float); + float *h_a[NSTREAMS], *h_b[NSTREAMS], *h_c[NSTREAMS]; + float *d_a[NSTREAMS], *d_b[NSTREAMS], *d_c[NSTREAMS]; + cudaStream_t st[NSTREAMS]; + cudaEvent_t ev[NSTREAMS], t0, t1; + float ms = 0.0f; + int bad = 0, i, s; + 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); + + for (s = 0; s < NSTREAMS; s++) { + h_a[s] = (float *)malloc(bytes); + h_b[s] = (float *)malloc(bytes); + h_c[s] = (float *)malloc(bytes); + if (!h_a[s] || !h_b[s] || !h_c[s]) + return 1; + for (i = 0; i < n; i++) { + h_a[s][i] = (float)((i + s) % 1000) * 0.5f; + h_b[s][i] = (float)((i + 2 * s) % 997) * 0.25f; + } + CHECK(cudaMalloc(&d_a[s], bytes)); + CHECK(cudaMalloc(&d_b[s], bytes)); + CHECK(cudaMalloc(&d_c[s], bytes)); + CHECK(cudaStreamCreate(&st[s])); + CHECK(cudaEventCreate(&ev[s])); + } + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + printf("DPTRS"); + for (s = 0; s < NSTREAMS; s++) + printf(" da%d=%p db%d=%p dc%d=%p", s, (void *)d_a[s], s, + (void *)d_b[s], s, (void *)d_c[s]); + printf("\n"); + printf("HPTRS"); + for (s = 0; s < NSTREAMS; s++) + printf(" ha%d=%p hb%d=%p hc%d=%p", s, (void *)h_a[s], s, + (void *)h_b[s], s, (void *)h_c[s]); + printf("\n"); + + CHECK(cudaEventRecord(t0)); + + for (s = 0; s < NSTREAMS; s++) { + CHECK(cudaMemcpyAsync(d_a[s], h_a[s], bytes, + cudaMemcpyHostToDevice, st[s])); + CHECK(cudaMemcpyAsync(d_b[s], h_b[s], bytes, + cudaMemcpyHostToDevice, st[s])); + vadd<<<grid, 256, 0, st[s]>>>(d_a[s], d_b[s], d_c[s], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[s], st[s])); + } + + CHECK(cudaStreamWaitEvent(st[1], ev[0], 0)); + vadd_into<<<grid, 256, 0, st[1]>>>(d_c[1], d_c[0], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[1], st[1])); + + CHECK(cudaStreamWaitEvent(st[2], ev[1], 0)); + vadd_into<<<grid, 256, 0, st[2]>>>(d_c[2], d_c[1], n); + CHECK(cudaGetLastError()); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaMemcpyAsync(h_c[s], d_c[s], bytes, + cudaMemcpyDeviceToHost, st[s])); + CHECK(cudaEventRecord(t1, st[2])); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaStreamSynchronize(st[s])); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + for (s = 0; s < NSTREAMS; s++) { + int b = 0; + float exp0[8], got0[8]; + + for (i = 0; i < n; i++) { + float c0 = h_a[0][i] + h_b[0][i]; + float c1 = h_a[1][i] + h_b[1][i] + c0; + float c2 = h_a[2][i] + h_b[2][i] + c1; + float want = (s == 0) ? c0 : (s == 1) ? c1 : c2; + + if (i < 8) { + exp0[i] = want; + got0[i] = h_c[s][i]; + } + if (h_c[s][i] != want) + b++; + } + bad += b; + printf("BUF s=%d bad=%d\n", s, b); + printf("FIRST s=%d exp=", s); + for (i = 0; i < 8; i++) + printf("%.4f ", exp0[i]); + printf(" got="); + for (i = 0; i < 8; i++) + printf("%.4f ", got0[i]); + printf("\n"); + } + + printf("STREAMS3 n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/tgemm_probe.cu b/tools/cuda-ab/tgemm_probe.cu new file mode 100644 index 0000000..f0a1084 --- /dev/null +++ b/tools/cuda-ab/tgemm_probe.cu @@ -0,0 +1,216 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * tgemm_probe.cu - minimal isolation for the transformer K/V divergence. + * + * Reproduces exactly the transformer's Q/K/V GEMM shape (X[64x256] x + * W[256x256], tile 16) for three weight matrices uploaded back-to-back + * (wq, wk, wv), then verifies each against a CPU reference and prints + * per-GEMM bad counts plus a bad-position histogram (first bad index per + * 16-row band and per 16-col band) so a copy-side hole is distinguishable + * from a kernel-side error. + * + * Output: + * DEVICE ... + * GEMM name=<q|k|v> bad=<n> first=(row,col) got=.. ref=.. + * BADPOS name=.. rows=<r0,r1,..> cols=<c0,c1,..> + * STATUS OK|FAIL + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cuda_runtime.h> + +#define SEQ 64 +#define DH 256 +#define TILE 16 + +__global__ void gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(t * TILE + threadIdx.y) * N + col]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[k][threadIdx.x]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +static float wval(int i, int seed) +{ + return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f; +} + +static float xval(int i, int j) +{ + return (float)(((i * 7 + j * 3 + 1) % 9) - 4) * 0.25f; +} + +int main(void) +{ + cudaDeviceProp prop; + float *h_x, *h_w, *h_c, *h_ref; + float *d_x, *d_w[3], *d_c[3]; + const int seqdh = SEQ * DH, dh2 = DH * DH; + int g, i, j, k; + + setvbuf(stdout, NULL, _IONBF, 0); + if (cudaGetDeviceProperties(&prop, 0) != cudaSuccess) { + printf("STATUS FAIL\n"); + return 1; + } + printf("DEVICE %s cc=%d.%d\n", prop.name, prop.major, prop.minor); + + h_x = (float *)malloc(seqdh * 4); + h_w = (float *)malloc(dh2 * 4); + h_c = (float *)malloc(seqdh * 4); + h_ref = (float *)malloc(seqdh * 4); + + if (cudaMalloc(&d_x, seqdh * 4) != cudaSuccess) + return 1; + for (g = 0; g < 3; g++) { + if (cudaMalloc(&d_w[g], dh2 * 4) != cudaSuccess) + return 1; + if (cudaMalloc(&d_c[g], seqdh * 4) != cudaSuccess) + return 1; + } + + for (i = 0; i < seqdh; i++) + h_x[i] = xval(i / DH, i % DH); + + /* Upload X once (the transformer's shared input). */ + if (cudaMemcpy(d_x, h_x, seqdh * 4, cudaMemcpyHostToDevice) != + cudaSuccess) + return 1; + /* Verify the X upload itself (first row + full buffer). */ + { + float *h_xb = (float *)malloc(seqdh * 4); + int xbad = 0, xr0 = 0; + + if (!h_xb) + return 1; + if (cudaMemcpy(h_xb, d_x, seqdh * 4, + cudaMemcpyDeviceToHost) != cudaSuccess) + return 1; + for (i = 0; i < seqdh; i++) + if (h_xb[i] != h_x[i]) { + xbad++; + if (i < DH) + xr0++; + } + printf("XVERIFY bad=%d row0_bad=%d\n", xbad, xr0); + free(h_xb); + } + + for (g = 0; g < 3; g++) { + int bad = 0, first_r = -1, first_c = -1; + float got = 0, ref = 0; + int rows[64] = {0}, cols[64] = {0}; + int nr = 0, nc = 0; + int pass; + + for (i = 0; i < dh2; i++) + h_w[i] = wval(i, g + 1); + if (cudaMemcpy(d_w[g], h_w, dh2 * 4, + cudaMemcpyHostToDevice) != cudaSuccess) + return 1; + + dim3 blk(TILE, TILE); + dim3 grd(DH / TILE, SEQ / TILE); + + for (pass = 0; pass < 3; pass++) { + float *hc = h_c, *hxb = h_ref; + int xbad = 0; + + gemm_nn<<<grd, blk>>>(d_x, d_w[g], d_c[g], SEQ, DH, DH); + if (cudaDeviceSynchronize() != cudaSuccess) { + printf("SYNC FAIL g=%d\n", g); + return 1; + } + if (cudaMemcpy(hc, d_c[g], seqdh * 4, + cudaMemcpyDeviceToHost) != cudaSuccess) + return 1; + /* re-verify X after the launch */ + if (cudaMemcpy(hxb, d_x, seqdh * 4, + cudaMemcpyDeviceToHost) != cudaSuccess) + return 1; + for (i = 0; i < seqdh; i++) + if (hxb[i] != h_x[i]) { + printf("XDIFF pass=%d i=%d got=%.6f want=%.6f\n", + pass, i, hxb[i], h_x[i]); + xbad++; + } + bad = 0; + nr = 0; + nc = 0; + for (i = 0; i < SEQ; i++) + for (j = 0; j < DH; j++) { + double acc = 0.0; + float r, a, diff; + + for (k = 0; k < DH; k++) + acc += (double)h_x[i * DH + k] * + (double)h_w[k * DH + j]; + r = (float)acc; + a = hc[i * DH + j]; + diff = a - r; + if (diff < 0) + diff = -diff; + if (diff > 1e-4f) { + int rr = i, cc = j; + + if (bad == 0) { + first_r = rr; + first_c = cc; + got = a; + ref = r; + } + bad++; + if (nr < 64 && (bad == 1 || + rr != rows[nr - 1])) + rows[nr++] = rr; + if (nc < 64 && (bad == 1 || + cc != cols[nc - 1])) + cols[nc++] = cc; + } + } + printf("PASS name=%c pass=%d bad=%d xbad=%d first=(%d,%d) got=%.6f ref=%.6f\n", + "qkv"[g], pass, bad, xbad, first_r, first_c, + got, ref); + if (bad) { + printf("ROW0 name=%c got=", "qkv"[g]); + for (i = 0; i < 8; i++) + printf("%.6f ", hc[i]); + printf(" ref="); + for (i = 0; i < 8; i++) { + double acc = 0.0; + + for (k = 0; k < DH; k++) + acc += (double)h_x[k] * + (double)h_w[k * DH + i]; + printf("%.6f ", (float)acc); + } + printf("\n"); + } + } + (void)got; + (void)ref; + (void)first_r; + (void)first_c; + (void)nr; + (void)nc; + } + + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/transformer.cu b/tools/cuda-ab/transformer.cu new file mode 100644 index 0000000..3044522 --- /dev/null +++ b/tools/cuda-ab/transformer.cu @@ -0,0 +1,472 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer. + * + * Two independent checks in one binary, both verified against an internal + * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor): + * + * 1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256) + * Q = X*Wq, K = X*Wk, V = X*Wv + * S = Q*K^T / 16 (scale = 1/sqrt(head_dim)) + * A = softmax(S) (row-wise) + * Y = (A*V)*Wo + * 2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then + * logits = h*Wh, P = softmax(logits) (row-wise over 256) + * + * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in + * fp32; only the softmax exponentials differ from the CPU reference, hence + * the relative tolerance. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum> + * FFN4 layers=4 bad=<mismatches> psum=<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 -arch=sm_75 -o transformer transformer.cu + * + * Then, from the guest: /tmp/transformer + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cmath> +#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 SEQ 64 +#define DH 256 +#define TILE 16 +#define TOL 1e-3f + +__global__ void gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(t * TILE + threadIdx.y) * N + col]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[k][threadIdx.x]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + /* C[i][j] = sum_k A[i][k] * B[j][k] (B is [N][K], row-major). */ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(blockIdx.x * TILE + threadIdx.y) * K + + t * TILE + threadIdx.x]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[threadIdx.x][k]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void softmax_rows(const float *S, float *P, int rows, int cols, + float scale) +{ + extern __shared__ float sm[]; + int row = blockIdx.x; + int t = threadIdx.x; + float v = S[row * cols + t] * scale; + int stride; + + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] = fmaxf(sm[t], sm[t + stride]); + __syncthreads(); + } + v = expf(v - sm[0]); + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] += sm[t + stride]; + __syncthreads(); + } + P[row * cols + t] = v / sm[0]; +} + +__global__ void relu_inplace(float *h, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n && h[i] < 0.0f) + h[i] = 0.0f; +} + +/* ------------------------------------------------------------------ host */ + +static void launch_gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_softmax(const float *S, float *P, int rows, int cols, + float scale) +{ + softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows, + cols, + scale); +} + +/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */ +static float xval(int i) +{ + return (float)(((i % 7) - 3)) * 0.25f; +} + +static float wval(int i, int seed) +{ + return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f; +} + +static void ref_gemm(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[k * N + j]; + C[i * N + j] = acc; + } +} + +static void ref_gemm_nt(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[j * K + k]; + C[i * N + j] = acc; + } +} + +static void ref_softmax(const double *S, double *P, int rows, int cols, + double scale) +{ + for (int i = 0; i < rows; i++) { + const double *s = S + i * cols; + double *p = P + i * cols; + double m = s[0] * scale, sum = 0.0; + + for (int j = 1; j < cols; j++) + if (s[j] * scale > m) + m = s[j] * scale; + for (int j = 0; j < cols; j++) { + p[j] = exp(s[j] * scale - m); + sum += p[j]; + } + for (int j = 0; j < cols; j++) + p[j] /= sum; + } +} + +static void ref_relu(double *h, int n) +{ + for (int i = 0; i < n; i++) + if (h[i] < 0.0) + h[i] = 0.0; +} + +static int count_bad(const float *gpu, const double *ref, int n) +{ + int bad = 0; + + for (int i = 0; i < n; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) + bad++; + } + return bad; +} + +static double dsum(const float *h, int n) +{ + double s = 0.0; + + for (int i = 0; i < n; i++) + s += (double)h[i]; + return s; +} + +#define NF 4 /* FFN layers */ + +int main(void) +{ + cudaDeviceProp prop; + const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ; + const size_t b_seqdh = (size_t)seqdh * sizeof(float); + const size_t b_dh2 = (size_t)dh2 * sizeof(float); + const size_t b_seq2 = (size_t)seq2 * sizeof(float); + float *h_x, *h_y, *h_p; + double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y; + double *r_f0, *r_f1, *r_logits, *r_p; + float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y; + float *d_wq, *d_wk, *d_wv, *d_wo; + float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p; + float *d_ping, *d_pong, *tmpf; + double *r_ping, *r_pong, *tmpd; + int bad_y, bad_p, i, l; + + 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); + + /* host inputs + CPU reference workspace */ + h_x = (float *)malloc(b_seqdh); + h_y = (float *)malloc(b_seqdh); + h_p = (float *)malloc(b_seqdh); + r_x = (double *)malloc((size_t)seqdh * sizeof(double)); + r_q = (double *)malloc((size_t)seqdh * sizeof(double)); + r_k = (double *)malloc((size_t)seqdh * sizeof(double)); + r_v = (double *)malloc((size_t)seqdh * sizeof(double)); + r_s = (double *)malloc((size_t)seq2 * sizeof(double)); + r_a = (double *)malloc((size_t)seq2 * sizeof(double)); + r_av = (double *)malloc((size_t)seqdh * sizeof(double)); + r_y = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f0 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f1 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_logits = (double *)malloc((size_t)seqdh * sizeof(double)); + r_p = (double *)malloc((size_t)seqdh * sizeof(double)); + if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s || + !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p) + return 1; + + for (i = 0; i < seqdh; i++) { + h_x[i] = xval(i); + r_x[i] = (double)xval(i); + } + + /* device buffers */ + CHECK(cudaMalloc(&d_x, b_seqdh)); + CHECK(cudaMalloc(&d_q, b_seqdh)); + CHECK(cudaMalloc(&d_k, b_seqdh)); + CHECK(cudaMalloc(&d_v, b_seqdh)); + CHECK(cudaMalloc(&d_s, b_seq2)); + CHECK(cudaMalloc(&d_a, b_seq2)); + CHECK(cudaMalloc(&d_av, b_seqdh)); + CHECK(cudaMalloc(&d_y, b_seqdh)); + CHECK(cudaMalloc(&d_wq, b_dh2)); + CHECK(cudaMalloc(&d_wk, b_dh2)); + CHECK(cudaMalloc(&d_wv, b_dh2)); + CHECK(cudaMalloc(&d_wo, b_dh2)); + CHECK(cudaMalloc(&d_wh, b_dh2)); + CHECK(cudaMalloc(&d_f0, b_seqdh)); + CHECK(cudaMalloc(&d_f1, b_seqdh)); + CHECK(cudaMalloc(&d_logits, b_seqdh)); + CHECK(cudaMalloc(&d_p, b_seqdh)); + for (l = 0; l < NF; l++) + CHECK(cudaMalloc(&d_wf[l], b_dh2)); + + /* weights: deterministic, uploaded once */ + { + float *w = (float *)malloc(b_dh2); + + if (!w) + return 1; + for (i = 0; i < dh2; i++) + w[i] = wval(i, 1); + CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 2); + CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 3); + CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 4); + CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 20); + CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice)); + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + w[i] = wval(i, 10 + l); + CHECK(cudaMemcpy(d_wf[l], w, b_dh2, + cudaMemcpyHostToDevice)); + } + free(w); + } + CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice)); + + /* ---- check 1: attention forward ---- */ + launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH); + launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH); + launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH); + launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH); + launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f); + launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ); + launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost)); + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w, r_q, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w, r_k, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w, r_v, SEQ, DH, DH); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w, r_y, SEQ, DH, DH); + free(r_w); + } + + bad_y = count_bad(h_y, r_y, seqdh); + printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y, + dsum(h_y, seqdh)); + + /* ---- check 2: 4-layer FFN stack + softmax head ---- */ + CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice)); + d_ping = d_f0; + d_pong = d_f1; + for (l = 0; l < NF; l++) { + launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH); + relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh); + tmpf = d_ping; + d_ping = d_pong; + d_pong = tmpf; + } + launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH); + launch_softmax(d_logits, d_p, SEQ, DH, 1.0f); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost)); + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double)); + r_ping = r_f0; + r_pong = r_f1; + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 10 + l); + ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH); + ref_relu(r_pong, seqdh); + tmpd = r_ping; + r_ping = r_pong; + r_pong = tmpd; + } + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 20); + ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH); + ref_softmax(r_logits, r_p, SEQ, DH, 1.0); + free(r_w); + } + + bad_p = count_bad(h_p, r_p, seqdh); + printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh)); + + if (bad_y != 0 || bad_p != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_x)); + CHECK(cudaFree(d_q)); + CHECK(cudaFree(d_k)); + CHECK(cudaFree(d_v)); + CHECK(cudaFree(d_s)); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_av)); + CHECK(cudaFree(d_y)); + CHECK(cudaFree(d_wq)); + CHECK(cudaFree(d_wk)); + CHECK(cudaFree(d_wv)); + CHECK(cudaFree(d_wo)); + CHECK(cudaFree(d_wh)); + for (l = 0; l < NF; l++) + CHECK(cudaFree(d_wf[l])); + CHECK(cudaFree(d_f0)); + CHECK(cudaFree(d_f1)); + CHECK(cudaFree(d_logits)); + CHECK(cudaFree(d_p)); + free(h_x); + free(h_y); + free(h_p); + free(r_x); + free(r_q); + free(r_k); + free(r_v); + free(r_s); + free(r_a); + free(r_av); + free(r_y); + free(r_f0); + free(r_f1); + free(r_logits); + free(r_p); + return 0; +} diff --git a/tools/cuda-ab/transformer_async.cu b/tools/cuda-ab/transformer_async.cu new file mode 100644 index 0000000..213363a --- /dev/null +++ b/tools/cuda-ab/transformer_async.cu @@ -0,0 +1,522 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer. + * + * Two independent checks in one binary, both verified against an internal + * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor): + * + * 1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256) + * Q = X*Wq, K = X*Wk, V = X*Wv + * S = Q*K^T / 16 (scale = 1/sqrt(head_dim)) + * A = softmax(S) (row-wise) + * Y = (A*V)*Wo + * 2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then + * logits = h*Wh, P = softmax(logits) (row-wise over 256) + * + * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in + * fp32; only the softmax exponentials differ from the CPU reference, hence + * the relative tolerance. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum> + * FFN4 layers=4 bad=<mismatches> psum=<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 -arch=sm_75 -o transformer transformer.cu + * + * Then, from the guest: /tmp/transformer + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cmath> +#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 SEQ 64 +#define DH 256 +#define TILE 16 +#define TOL 1e-3f + +__global__ void gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(t * TILE + threadIdx.y) * N + col]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[k][threadIdx.x]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + /* C[i][j] = sum_k A[i][k] * B[j][k] (B is [N][K], row-major). */ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(blockIdx.x * TILE + threadIdx.y) * K + + t * TILE + threadIdx.x]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[threadIdx.x][k]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void softmax_rows(const float *S, float *P, int rows, int cols, + float scale) +{ + extern __shared__ float sm[]; + int row = blockIdx.x; + int t = threadIdx.x; + float v = S[row * cols + t] * scale; + int stride; + + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] = fmaxf(sm[t], sm[t + stride]); + __syncthreads(); + } + v = expf(v - sm[0]); + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] += sm[t + stride]; + __syncthreads(); + } + P[row * cols + t] = v / sm[0]; +} + +__global__ void relu_inplace(float *h, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n && h[i] < 0.0f) + h[i] = 0.0f; +} + +/* ------------------------------------------------------------------ host */ + +static void launch_gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_softmax(const float *S, float *P, int rows, int cols, + float scale) +{ + softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows, + cols, + scale); +} + +/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */ +static float xval(int i) +{ + return (float)(((i % 7) - 3)) * 0.25f; +} + +static float wval(int i, int seed) +{ + return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f; +} + +static void ref_gemm(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[k * N + j]; + C[i * N + j] = acc; + } +} + +static void ref_gemm_nt(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[j * K + k]; + C[i * N + j] = acc; + } +} + +static void ref_softmax(const double *S, double *P, int rows, int cols, + double scale) +{ + for (int i = 0; i < rows; i++) { + const double *s = S + i * cols; + double *p = P + i * cols; + double m = s[0] * scale, sum = 0.0; + + for (int j = 1; j < cols; j++) + if (s[j] * scale > m) + m = s[j] * scale; + for (int j = 0; j < cols; j++) { + p[j] = exp(s[j] * scale - m); + sum += p[j]; + } + for (int j = 0; j < cols; j++) + p[j] /= sum; + } +} + +static void ref_relu(double *h, int n) +{ + for (int i = 0; i < n; i++) + if (h[i] < 0.0) + h[i] = 0.0; +} + +static int count_bad(const float *gpu, const double *ref, int n) +{ + int bad = 0; + + for (int i = 0; i < n; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) + bad++; + } + return bad; +} + +static double dsum(const float *h, int n) +{ + double s = 0.0; + + for (int i = 0; i < n; i++) + s += (double)h[i]; + return s; +} + +#define NF 4 /* FFN layers */ + +int main(void) +{ + cudaDeviceProp prop; + const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ; + const size_t b_seqdh = (size_t)seqdh * sizeof(float); + const size_t b_dh2 = (size_t)dh2 * sizeof(float); + const size_t b_seq2 = (size_t)seq2 * sizeof(float); + float *h_x, *h_y, *h_p; + double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y; + double *r_f0, *r_f1, *r_logits, *r_p; + float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y; + float *d_wq, *d_wk, *d_wv, *d_wo; + float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p; + float *d_ping, *d_pong, *tmpf; + double *r_ping, *r_pong, *tmpd; + int bad_y, bad_p, i, l; + + 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); + + /* host inputs + CPU reference workspace */ + h_x = (float *)malloc(b_seqdh); + h_y = (float *)malloc(b_seqdh); + h_p = (float *)malloc(b_seqdh); + r_x = (double *)malloc((size_t)seqdh * sizeof(double)); + r_q = (double *)malloc((size_t)seqdh * sizeof(double)); + r_k = (double *)malloc((size_t)seqdh * sizeof(double)); + r_v = (double *)malloc((size_t)seqdh * sizeof(double)); + r_s = (double *)malloc((size_t)seq2 * sizeof(double)); + r_a = (double *)malloc((size_t)seq2 * sizeof(double)); + r_av = (double *)malloc((size_t)seqdh * sizeof(double)); + r_y = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f0 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f1 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_logits = (double *)malloc((size_t)seqdh * sizeof(double)); + r_p = (double *)malloc((size_t)seqdh * sizeof(double)); + if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s || + !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p) + return 1; + + for (i = 0; i < seqdh; i++) { + h_x[i] = xval(i); + r_x[i] = (double)xval(i); + } + + /* device buffers */ + CHECK(cudaMalloc(&d_x, b_seqdh)); + CHECK(cudaMalloc(&d_q, b_seqdh)); + CHECK(cudaMalloc(&d_k, b_seqdh)); + CHECK(cudaMalloc(&d_v, b_seqdh)); + CHECK(cudaMalloc(&d_s, b_seq2)); + CHECK(cudaMalloc(&d_a, b_seq2)); + CHECK(cudaMalloc(&d_av, b_seqdh)); + CHECK(cudaMalloc(&d_y, b_seqdh)); + CHECK(cudaMalloc(&d_wq, b_dh2)); + CHECK(cudaMalloc(&d_wk, b_dh2)); + CHECK(cudaMalloc(&d_wv, b_dh2)); + CHECK(cudaMalloc(&d_wo, b_dh2)); + CHECK(cudaMalloc(&d_wh, b_dh2)); + CHECK(cudaMalloc(&d_f0, b_seqdh)); + CHECK(cudaMalloc(&d_f1, b_seqdh)); + CHECK(cudaMalloc(&d_logits, b_seqdh)); + CHECK(cudaMalloc(&d_p, b_seqdh)); + for (l = 0; l < NF; l++) + CHECK(cudaMalloc(&d_wf[l], b_dh2)); + + /* weights: deterministic, uploaded once */ + { + float *w = (float *)malloc(b_dh2); + + if (!w) + return 1; + for (i = 0; i < dh2; i++) + w[i] = wval(i, 1); + CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 2); + CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 3); + CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 4); + CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 20); + CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice)); + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + w[i] = wval(i, 10 + l); + CHECK(cudaMemcpy(d_wf[l], w, b_dh2, + cudaMemcpyHostToDevice)); + } + free(w); + } + CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice)); + + /* ---- check 1: attention forward (async chain, stage readbacks) ---- */ + launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH); + launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH); + launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH); + launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH); + launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f); + launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ); + launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost)); + { + float *hs = (float *)malloc(b_seqdh); + float *hs2 = (float *)malloc(b_seq2); + double *rw = (double *)malloc((size_t)dh2 * sizeof(double)); + double *rq = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rk = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rv = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rs = (double *)malloc((size_t)seq2 * sizeof(double)); + double *ra = (double *)malloc((size_t)seq2 * sizeof(double)); + double *rav = (double *)malloc((size_t)seqdh * sizeof(double)); + double *ry = (double *)malloc((size_t)seqdh * sizeof(double)); + int st; + + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 1); + ref_gemm(r_x, rw, rq, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 2); + ref_gemm(r_x, rw, rk, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 3); + ref_gemm(r_x, rw, rv, SEQ, DH, DH); + ref_gemm_nt(rq, rk, rs, SEQ, SEQ, DH); + ref_softmax(rs, ra, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(ra, rv, rav, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 4); + ref_gemm(rav, rw, ry, SEQ, DH, DH); + + CHECK(cudaMemcpy(hs, d_q, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rq, seqdh); + printf("ASYNC Q bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_k, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rk, seqdh); + printf("ASYNC K bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_v, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rv, seqdh); + printf("ASYNC V bad=%d\n", st); + CHECK(cudaMemcpy(hs2, d_s, b_seq2, cudaMemcpyDeviceToHost)); + st = count_bad(hs2, rs, seq2); + printf("ASYNC S bad=%d\n", st); + CHECK(cudaMemcpy(hs2, d_a, b_seq2, cudaMemcpyDeviceToHost)); + st = count_bad(hs2, ra, seq2); + printf("ASYNC A bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_av, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rav, seqdh); + printf("ASYNC AV bad=%d\n", st); + st = count_bad(h_y, ry, seqdh); + printf("ASYNC Y bad=%d\n", st); + } + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w, r_q, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w, r_k, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w, r_v, SEQ, DH, DH); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w, r_y, SEQ, DH, DH); + free(r_w); + } + + bad_y = count_bad(h_y, r_y, seqdh); + printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y, + dsum(h_y, seqdh)); + + /* ---- check 2: 4-layer FFN stack + softmax head ---- */ + CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice)); + d_ping = d_f0; + d_pong = d_f1; + for (l = 0; l < NF; l++) { + launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH); + relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh); + tmpf = d_ping; + d_ping = d_pong; + d_pong = tmpf; + } + launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH); + launch_softmax(d_logits, d_p, SEQ, DH, 1.0f); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost)); + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double)); + r_ping = r_f0; + r_pong = r_f1; + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 10 + l); + ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH); + ref_relu(r_pong, seqdh); + tmpd = r_ping; + r_ping = r_pong; + r_pong = tmpd; + } + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 20); + ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH); + ref_softmax(r_logits, r_p, SEQ, DH, 1.0); + free(r_w); + } + + bad_p = count_bad(h_p, r_p, seqdh); + printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh)); + + if (bad_y != 0 || bad_p != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_x)); + CHECK(cudaFree(d_q)); + CHECK(cudaFree(d_k)); + CHECK(cudaFree(d_v)); + CHECK(cudaFree(d_s)); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_av)); + CHECK(cudaFree(d_y)); + CHECK(cudaFree(d_wq)); + CHECK(cudaFree(d_wk)); + CHECK(cudaFree(d_wv)); + CHECK(cudaFree(d_wo)); + CHECK(cudaFree(d_wh)); + for (l = 0; l < NF; l++) + CHECK(cudaFree(d_wf[l])); + CHECK(cudaFree(d_f0)); + CHECK(cudaFree(d_f1)); + CHECK(cudaFree(d_logits)); + CHECK(cudaFree(d_p)); + free(h_x); + free(h_y); + free(h_p); + free(r_x); + free(r_q); + free(r_k); + free(r_v); + free(r_s); + free(r_a); + free(r_av); + free(r_y); + free(r_f0); + free(r_f1); + free(r_logits); + free(r_p); + return 0; +} diff --git a/tools/cuda-ab/transformer_stages.cu b/tools/cuda-ab/transformer_stages.cu new file mode 100644 index 0000000..d7f67c0 --- /dev/null +++ b/tools/cuda-ab/transformer_stages.cu @@ -0,0 +1,553 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer. + * + * Two independent checks in one binary, both verified against an internal + * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor): + * + * 1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256) + * Q = X*Wq, K = X*Wk, V = X*Wv + * S = Q*K^T / 16 (scale = 1/sqrt(head_dim)) + * A = softmax(S) (row-wise) + * Y = (A*V)*Wo + * 2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then + * logits = h*Wh, P = softmax(logits) (row-wise over 256) + * + * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in + * fp32; only the softmax exponentials differ from the CPU reference, hence + * the relative tolerance. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum> + * FFN4 layers=4 bad=<mismatches> psum=<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 -arch=sm_75 -o transformer transformer.cu + * + * Then, from the guest: /tmp/transformer + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cmath> +#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 SEQ 64 +#define DH 256 +#define TILE 16 +#define TOL 1e-3f + +__global__ void gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(t * TILE + threadIdx.y) * N + col]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[k][threadIdx.x]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + /* C[i][j] = sum_k A[i][k] * B[j][k] (B is [N][K], row-major). */ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(blockIdx.x * TILE + threadIdx.y) * K + + t * TILE + threadIdx.x]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[threadIdx.x][k]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void softmax_rows(const float *S, float *P, int rows, int cols, + float scale) +{ + extern __shared__ float sm[]; + int row = blockIdx.x; + int t = threadIdx.x; + float v = S[row * cols + t] * scale; + int stride; + + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] = fmaxf(sm[t], sm[t + stride]); + __syncthreads(); + } + v = expf(v - sm[0]); + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] += sm[t + stride]; + __syncthreads(); + } + P[row * cols + t] = v / sm[0]; +} + +__global__ void relu_inplace(float *h, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n && h[i] < 0.0f) + h[i] = 0.0f; +} + +/* ------------------------------------------------------------------ host */ + +static void launch_gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_softmax(const float *S, float *P, int rows, int cols, + float scale) +{ + softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows, + cols, + scale); +} + +/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */ +static float xval(int i) +{ + return (float)(((i % 7) - 3)) * 0.25f; +} + +static float wval(int i, int seed) +{ + return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f; +} + +static void ref_gemm(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[k * N + j]; + C[i * N + j] = acc; + } +} + +static void ref_gemm_nt(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[j * K + k]; + C[i * N + j] = acc; + } +} + +static void ref_softmax(const double *S, double *P, int rows, int cols, + double scale) +{ + for (int i = 0; i < rows; i++) { + const double *s = S + i * cols; + double *p = P + i * cols; + double m = s[0] * scale, sum = 0.0; + + for (int j = 1; j < cols; j++) + if (s[j] * scale > m) + m = s[j] * scale; + for (int j = 0; j < cols; j++) { + p[j] = exp(s[j] * scale - m); + sum += p[j]; + } + for (int j = 0; j < cols; j++) + p[j] /= sum; + } +} + +static void ref_relu(double *h, int n) +{ + for (int i = 0; i < n; i++) + if (h[i] < 0.0) + h[i] = 0.0; +} + +static int count_bad(const float *gpu, const double *ref, int n) +{ + int bad = 0; + + for (int i = 0; i < n; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) + bad++; + } + return bad; +} + +static double dsum(const float *h, int n) +{ + double s = 0.0; + + for (int i = 0; i < n; i++) + s += (double)h[i]; + return s; +} + +static void dump_bad(const char *tag, const float *gpu, const double *ref, + int n) +{ + int i, shown = 0; + + for (i = 0; i < n && shown < 8; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) { + printf("BAD %s i=%d gpu=%.6f ref=%.6f (tile %d,%d row %d col %d)\n", + tag, i, g, r, i / 16, (i % 256) / 16, + i / 256, i % 256); + shown++; + } + } +} + +#define NF 4 /* FFN layers */ + +int main(void) +{ + cudaDeviceProp prop; + const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ; + const size_t b_seqdh = (size_t)seqdh * sizeof(float); + const size_t b_dh2 = (size_t)dh2 * sizeof(float); + const size_t b_seq2 = (size_t)seq2 * sizeof(float); + float *h_x, *h_y, *h_p; + double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y; + double *r_f0, *r_f1, *r_logits, *r_p; + float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y; + float *d_wq, *d_wk, *d_wv, *d_wo; + float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p; + float *d_ping, *d_pong, *tmpf; + double *r_ping, *r_pong, *tmpd; + int bad_y, bad_p, i, l; + + 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); + + /* host inputs + CPU reference workspace */ + h_x = (float *)malloc(b_seqdh); + h_y = (float *)malloc(b_seqdh); + h_p = (float *)malloc(b_seqdh); + r_x = (double *)malloc((size_t)seqdh * sizeof(double)); + r_q = (double *)malloc((size_t)seqdh * sizeof(double)); + r_k = (double *)malloc((size_t)seqdh * sizeof(double)); + r_v = (double *)malloc((size_t)seqdh * sizeof(double)); + r_s = (double *)malloc((size_t)seq2 * sizeof(double)); + r_a = (double *)malloc((size_t)seq2 * sizeof(double)); + r_av = (double *)malloc((size_t)seqdh * sizeof(double)); + r_y = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f0 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f1 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_logits = (double *)malloc((size_t)seqdh * sizeof(double)); + r_p = (double *)malloc((size_t)seqdh * sizeof(double)); + if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s || + !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p) + return 1; + + for (i = 0; i < seqdh; i++) { + h_x[i] = xval(i); + r_x[i] = (double)xval(i); + } + + /* device buffers */ + CHECK(cudaMalloc(&d_x, b_seqdh)); + CHECK(cudaMalloc(&d_q, b_seqdh)); + CHECK(cudaMalloc(&d_k, b_seqdh)); + CHECK(cudaMalloc(&d_v, b_seqdh)); + CHECK(cudaMalloc(&d_s, b_seq2)); + CHECK(cudaMalloc(&d_a, b_seq2)); + CHECK(cudaMalloc(&d_av, b_seqdh)); + CHECK(cudaMalloc(&d_y, b_seqdh)); + CHECK(cudaMalloc(&d_wq, b_dh2)); + CHECK(cudaMalloc(&d_wk, b_dh2)); + CHECK(cudaMalloc(&d_wv, b_dh2)); + CHECK(cudaMalloc(&d_wo, b_dh2)); + CHECK(cudaMalloc(&d_wh, b_dh2)); + CHECK(cudaMalloc(&d_f0, b_seqdh)); + CHECK(cudaMalloc(&d_f1, b_seqdh)); + CHECK(cudaMalloc(&d_logits, b_seqdh)); + CHECK(cudaMalloc(&d_p, b_seqdh)); + for (l = 0; l < NF; l++) + CHECK(cudaMalloc(&d_wf[l], b_dh2)); + + /* weights: deterministic, uploaded once */ + { + float *w = (float *)malloc(b_dh2); + + if (!w) + return 1; + for (i = 0; i < dh2; i++) + w[i] = wval(i, 1); + CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 2); + CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 3); + CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 4); + CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 20); + CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice)); + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + w[i] = wval(i, 10 + l); + CHECK(cudaMemcpy(d_wf[l], w, b_dh2, + cudaMemcpyHostToDevice)); + } + free(w); + } + CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice)); + + /* ---- check 1: attention forward, stage-by-stage verify ---- */ + { + float *h_stage = (float *)malloc(b_seqdh); + float *h_s2 = (float *)malloc(b_seq2); + double *r_stage = (double *)malloc((size_t)seqdh * sizeof(double)); + double *r_s2 = (double *)malloc((size_t)seq2 * sizeof(double)); + double *r_w2 = (double *)malloc((size_t)dh2 * sizeof(double)); + int st; + + if (!h_stage || !h_s2 || !r_stage || !r_s2 || !r_w2) + return 1; + + launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_q, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE Q bad=%d\n", st); + dump_bad("Q", h_stage, r_stage, seqdh); + + launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_k, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE K bad=%d\n", st); + dump_bad("K", h_stage, r_stage, seqdh); + + launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_v, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE V bad=%d\n", st); + dump_bad("V", h_stage, r_stage, seqdh); + + launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_s2, d_s, b_seq2, cudaMemcpyDeviceToHost)); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + st = count_bad(h_s2, r_s, seq2); + printf("STAGE S bad=%d\n", st); + dump_bad("S", h_s2, r_s, seq2); + + launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_s2, d_a, b_seq2, cudaMemcpyDeviceToHost)); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + st = count_bad(h_s2, r_a, seq2); + printf("STAGE A bad=%d\n", st); + dump_bad("A", h_s2, r_a, seq2); + + launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_av, b_seqdh, cudaMemcpyDeviceToHost)); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + st = count_bad(h_stage, r_av, seqdh); + printf("STAGE AV bad=%d\n", st); + + launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w2, r_y, SEQ, DH, DH); + st = count_bad(h_y, r_y, seqdh); + printf("STAGE Y bad=%d\n", st); + } + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w, r_q, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w, r_k, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w, r_v, SEQ, DH, DH); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w, r_y, SEQ, DH, DH); + free(r_w); + } + + bad_y = count_bad(h_y, r_y, seqdh); + printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y, + dsum(h_y, seqdh)); + + /* ---- check 2: 4-layer FFN stack + softmax head ---- */ + CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice)); + d_ping = d_f0; + d_pong = d_f1; + for (l = 0; l < NF; l++) { + launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH); + relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh); + tmpf = d_ping; + d_ping = d_pong; + d_pong = tmpf; + } + launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH); + launch_softmax(d_logits, d_p, SEQ, DH, 1.0f); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost)); + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double)); + r_ping = r_f0; + r_pong = r_f1; + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 10 + l); + ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH); + ref_relu(r_pong, seqdh); + tmpd = r_ping; + r_ping = r_pong; + r_pong = tmpd; + } + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 20); + ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH); + ref_softmax(r_logits, r_p, SEQ, DH, 1.0); + free(r_w); + } + + bad_p = count_bad(h_p, r_p, seqdh); + printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh)); + + if (bad_y != 0 || bad_p != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_x)); + CHECK(cudaFree(d_q)); + CHECK(cudaFree(d_k)); + CHECK(cudaFree(d_v)); + CHECK(cudaFree(d_s)); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_av)); + CHECK(cudaFree(d_y)); + CHECK(cudaFree(d_wq)); + CHECK(cudaFree(d_wk)); + CHECK(cudaFree(d_wv)); + CHECK(cudaFree(d_wo)); + CHECK(cudaFree(d_wh)); + for (l = 0; l < NF; l++) + CHECK(cudaFree(d_wf[l])); + CHECK(cudaFree(d_f0)); + CHECK(cudaFree(d_f1)); + CHECK(cudaFree(d_logits)); + CHECK(cudaFree(d_p)); + free(h_x); + free(h_y); + free(h_p); + free(r_x); + free(r_q); + free(r_k); + free(r_v); + free(r_s); + free(r_a); + free(r_av); + free(r_y); + free(r_f0); + free(r_f1); + free(r_logits); + free(r_p); + return 0; +} diff --git a/tools/cuda-ab/wedge_ab.cu b/tools/cuda-ab/wedge_ab.cu new file mode 100644 index 0000000..2d88f92 --- /dev/null +++ b/tools/cuda-ab/wedge_ab.cu @@ -0,0 +1,100 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_ab.cu - discriminate the wedge trigger: is it the second launch of + * the SAME program, or the second launch of the context regardless of + * program? + * + * Sequence (each launch verified against mutated inputs): + * 1. vadd (program A) + * 2. vsub (program B, different code) + * 3. vadd (program A again) + * 4. vsub (program B again) + * + * If launch 2 (B) executes -> the trigger is same-program reuse. + * If launch 2 (B) wedges -> the trigger is the second launch in the context. + * + * Build: nvcc -O2 --cudart=static -o wedge_ab wedge_ab.cu + */ + +#include <cstdio> +#include <cstdlib> +#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) + +__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]; +} + +__global__ void vsub(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]; +} + +int main(void) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc; + float *da, *db, *dc; + int i, bad, step; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + if (!ha || !hb || !hc) + return 1; + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + + for (step = 0; step < 4; step++) { + const char *which = (step & 1) ? "vsub" : "vadd"; + int is_sub = step & 1; + + if (step > 0) { + ha[step] += 1.0f; + hb[step * 7 + 3] += 2.0f; + } + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + if (is_sub) + vsub<<<(n + 255) / 256, 256>>>(da, db, dc, n); + else + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + bad = 0; + for (i = 0; i < n; i++) { + float expect = is_sub ? ha[i] - hb[i] : ha[i] + hb[i]; + + if (hc[i] != expect) + bad++; + } + printf("STEP %d %s bad=%d/%d\n", step, which, bad, n); + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/wedge_bisect.cu b/tools/cuda-ab/wedge_bisect.cu new file mode 100644 index 0000000..e062a2b --- /dev/null +++ b/tools/cuda-ab/wedge_bisect.cu @@ -0,0 +1,104 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_bisect.cu - bisect the wedge trigger: which inter-launch activity + * makes the SECOND vadd launch no-execute? + * + * Every iteration is: mutate host inputs -> H2D da/db -> launch -> sync -> + * D2H dc -> verify (the wedge_recover shape). Modes add activity between + * the verify and the next launch: + * + * mode A: none (pure second-launch class) + * mode B: an extra H2D of the same buffers (CE write before the launch) + * mode C: an extra D2H readback (CE read before the launch) + * mode D: a 500 ms sleep (no CE traffic) + * + * A no-execute launch is reported (bad>0) not masked: the inputs mutate every + * iteration. + * + * Build: nvcc -O2 --cudart=static -o wedge_bisect wedge_bisect.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#include <cuda_runtime.h> +#include <unistd.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) + +__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]; +} + +int main(int argc, char **argv) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc; + float *da, *db, *dc; + const char *mode = argc > 1 ? argv[1] : "A"; + int iters = argc > 2 ? atoi(argv[2]) : 4; + int i, bad = 0; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + if (!ha || !hb || !hc) + return 1; + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + for (i = 0; i < iters; i++) { + int j, itbad = 0; + + if (i > 0) { + ha[i % n] += 1.0f; + hb[(i * 7 + 3) % n] += 2.0f; + } + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + for (j = 0; j < n; j++) + if (hc[j] != ha[j] + hb[j]) + itbad++; + printf("MODE %s iter=%d bad=%d/%d\n", mode, i, itbad, n); + bad += itbad; + if (i + 1 < iters) { + if (mode[0] == 'B') { + CHECK(cudaMemcpy(da, ha, bytes, + cudaMemcpyHostToDevice)); + } else if (mode[0] == 'C') { + CHECK(cudaMemcpy(hc, dc, bytes, + cudaMemcpyDeviceToHost)); + } else if (mode[0] == 'D') { + usleep(500000); + } + } + } + printf("STATUS %s\n", bad ? "FAIL" : "OK"); + return bad ? 1 : 0; +} diff --git a/tools/cuda-ab/wedge_recover.cu b/tools/cuda-ab/wedge_recover.cu new file mode 100644 index 0000000..fd8f3f9 --- /dev/null +++ b/tools/cuda-ab/wedge_recover.cu @@ -0,0 +1,143 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_recover.cu - reproduce the per-launch no-execute (wedge) and test + * whether an idle gap heals a wedged context. + * + * v2 (session 10, r416 finding): the v1 loop reused one deterministic + * kernel+buffers, so a no-write launch was MASKED - the stale buffer held + * exactly the previous iteration's (identical) values and the verify passed. + * r416 ran 60/60 "OK" while 58 consecutive launches no-wrote underneath. + * + * v2 mutates the inputs every iteration (H2D of fresh ha/hb), so a stale + * readback cannot match the new expected values: a no-write launch must + * report BAD. On BAD the same iteration is retried after a sleep, and the + * readback is classified: STALE-PREV (byte-identical to the previous + * iteration's output = launch never executed) vs PARTIAL/OTHER (engine wrote + * something). Latency and per-iteration outcome are printed. + * + * Build: nvcc -O2 --cudart=static -o wedge_recover wedge_recover.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cuda_runtime.h> +#include <unistd.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) + +__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]; +} + +int main(int argc, char **argv) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc, *hc_prev; + float *da, *db, *dc; + int iters = argc > 1 ? atoi(argv[1]) : 40; + int sleep_s = argc > 2 ? atoi(argv[2]) : 20; + int fails = 0, heals = 0, i; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + hc_prev = (float *)malloc(bytes); + if (!ha || !hb || !hc || !hc_prev) + return 1; + /* iteration 0 baseline; mutated each iteration below (v2) */ + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + memset(hc_prev, 0, bytes); + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + for (i = 0; i < iters; i++) { + int attempt = 0; + long long ms[2] = { 0, 0 }; + + /* + * v2: fresh inputs every iteration so a stale (never + * executed) launch cannot match. Mutate after iteration 0. + */ + if (i > 0) { + ha[i % n] += 1.0f; + hb[(i * 7 + 3) % n] += 2.0f; + CHECK(cudaMemcpy(da, ha, bytes, + cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, + cudaMemcpyHostToDevice)); + } + + for (;;) { + int bad = 0, stale_prev = 1, j; + cudaEvent_t e0, e1; + float ems = 0; + + cudaEventCreate(&e0); + cudaEventCreate(&e1); + cudaEventRecord(e0); + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + cudaEventRecord(e1); + CHECK(cudaDeviceSynchronize()); + cudaEventElapsedTime(&ems, e0, e1); + cudaEventDestroy(e0); + cudaEventDestroy(e1); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + for (j = 0; j < n; j++) { + float expect = ha[j] + hb[j]; + + if (hc[j] != expect) { + bad++; + if (hc[j] != hc_prev[j]) + stale_prev = 0; + } + } + ms[attempt & 1] = (long long)(ems * 1000.0f); + if (bad == 0) { + if (attempt > 0) + heals++; + printf("ITER %d OK attempt=%d launch_us=%lld\n", + i, attempt, ms[attempt & 1]); + memcpy(hc_prev, hc, bytes); + break; + } + fails++; + printf("ITER %d BAD bad=%d/%d attempt=%d launch_us=%lld class=%s\n", + i, bad, n, attempt, ms[attempt & 1], + stale_prev ? "STALE-PREV(NOWRITE)" + : "OTHER/PARTIAL"); + if (attempt >= 3) { + printf("ITER %d GIVEUP\n", i); + break; + } + attempt++; + printf("ITER %d SLEEP %ds before retry\n", i, sleep_s); + sleep(sleep_s); + } + } + printf("SUMMARY iters=%d fails=%d heals=%d sleep_s=%d\n", + iters, fails, heals, sleep_s); + printf("STATUS %s\n", fails == 0 ? "OK" : "FAIL"); + return fails == 0 ? 0 : 1; +} diff --git a/tools/guest-diag/egl_frame.c b/tools/guest-diag/egl_frame.c new file mode 100644 index 0000000..ed10e39 --- /dev/null +++ b/tools/guest-diag/egl_frame.c @@ -0,0 +1,182 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * egl_frame.c - minimal EGL/GLES2 offscreen render test for the vGPU guest. + * + * Tries, in order: the EGL_EXT_platform_device platform with each enumerated + * EGLDeviceEXT, EGL_MESA_platform_surfaceless, then EGL_DEFAULT_DISPLAY. + * For the first display that initializes it creates an offscreen pbuffer + * context, renders colour bands into an FBO, reads them back and writes a + * PPM. One greppable line per stage. + * + * Build (guest): gcc -O2 -o egl_frame egl_frame.c -lEGL -lGLESv2 + */ +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include <EGL/egl.h> +#include <EGL/eglext.h> +#include <GLES2/gl2.h> + +#define W 320 +#define H 240 + +static EGLDisplay try_display(const char *name, EGLDisplay dpy, + const char *dev_name) +{ + EGLint major = 0, minor = 0; + + if (dpy == EGL_NO_DISPLAY) { + printf("EGL_PLATFORM %s -> no display\n", name); + return EGL_NO_DISPLAY; + } + if (!eglInitialize(dpy, &major, &minor)) { + printf("EGL_PLATFORM %s -> init fail 0x%x\n", name, + eglGetError()); + return EGL_NO_DISPLAY; + } + printf("EGL_PLATFORM %s ok %d.%d vendor=%s%s%s\n", name, major, minor, + eglQueryString(dpy, EGL_VENDOR), + dev_name ? " device=" : "", dev_name ? dev_name : ""); + return dpy; +} + +static EGLDisplay get_display(void) +{ + PFNEGLGETPLATFORMDISPLAYEXTPROC gp = + (PFNEGLGETPLATFORMDISPLAYEXTPROC)eglGetProcAddress( + "eglGetPlatformDisplayEXT"); + PFNEGLQUERYDEVICESEXTPROC qd = + (PFNEGLQUERYDEVICESEXTPROC)eglGetProcAddress( + "eglQueryDevicesEXT"); + EGLDisplay dpy; + + if (gp && qd) { + EGLDeviceEXT devs[8]; + EGLint n = 0; + int i; + + if (qd(8, devs, &n) && n > 0) { + printf("EGL_DEVICES n=%d\n", n); + for (i = 0; i < n; i++) { + EGLDisplay d = gp(EGL_PLATFORM_DEVICE_EXT, + devs[i], NULL); + + dpy = try_display("device", d, "enum"); + if (dpy != EGL_NO_DISPLAY) + return dpy; + } + } else { + printf("EGL_DEVICES n=0\n"); + } + if (gp) { + dpy = try_display("surfaceless", + gp(EGL_PLATFORM_SURFACELESS_MESA, + NULL, NULL), NULL); + if (dpy != EGL_NO_DISPLAY) + return dpy; + } + } + return try_display("default", eglGetDisplay(EGL_DEFAULT_DISPLAY), + NULL); +} + +static int render(EGLDisplay dpy) +{ + EGLConfig cfg; + EGLint ncfg = 0; + EGLSurface surf; + EGLContext ctx; + GLuint fbo = 0, tex = 0; + unsigned char *px; + FILE *f; + int i; + EGLint attrs[] = { + EGL_SURFACE_TYPE, EGL_PBUFFER_BIT, + EGL_RENDERABLE_TYPE, EGL_OPENGL_ES2_BIT, + EGL_RED_SIZE, 8, EGL_GREEN_SIZE, 8, + EGL_BLUE_SIZE, 8, EGL_ALPHA_SIZE, 8, + EGL_NONE + }; + EGLint pattrs[] = { EGL_WIDTH, W, EGL_HEIGHT, H, EGL_NONE }; + EGLint cattrs[] = { EGL_CONTEXT_CLIENT_VERSION, 2, EGL_NONE }; + + if (!eglChooseConfig(dpy, attrs, &cfg, 1, &ncfg) || ncfg < 1) { + printf("EGL_CONFIG_FAIL err=0x%x\n", eglGetError()); + return 1; + } + surf = eglCreatePbufferSurface(dpy, cfg, pattrs); + if (surf == EGL_NO_SURFACE) { + printf("EGL_PBUFFER_FAIL err=0x%x\n", eglGetError()); + return 1; + } + ctx = eglCreateContext(dpy, cfg, EGL_NO_CONTEXT, cattrs); + if (ctx == EGL_NO_CONTEXT) { + printf("EGL_CONTEXT_FAIL err=0x%x\n", eglGetError()); + return 1; + } + if (!eglMakeCurrent(dpy, surf, surf, ctx)) { + printf("EGL_CURRENT_FAIL err=0x%x\n", eglGetError()); + return 1; + } + printf("GL_RENDERER %s\n", glGetString(GL_RENDERER)); + printf("GL_VERSION %s\n", glGetString(GL_VERSION)); + + glGenTextures(1, &tex); + glBindTexture(GL_TEXTURE_2D, tex); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, W, H, 0, GL_RGBA, + GL_UNSIGNED_BYTE, NULL); + glGenFramebuffers(1, &fbo); + glBindFramebuffer(GL_FRAMEBUFFER, fbo); + glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, + GL_TEXTURE_2D, tex, 0); + if (glCheckFramebufferStatus(GL_FRAMEBUFFER) != + GL_FRAMEBUFFER_COMPLETE) { + printf("GL_FBO_FAIL\n"); + return 1; + } + for (i = 0; i < 16; i++) { + glEnable(GL_SCISSOR_TEST); + glScissor(0, i * H / 16, W, H / 16); + glClearColor((float)i / 16.0f, 1.0f - (float)i / 16.0f, + 0.5f, 1.0f); + glClear(GL_COLOR_BUFFER_BIT); + } + glFinish(); + printf("GL_ERR 0x%x\n", glGetError()); + px = (unsigned char *)malloc(W * H * 4); + glReadPixels(0, 0, W, H, GL_RGBA, GL_UNSIGNED_BYTE, px); + { + const char *out = getenv("EGL_FRAME_OUT"); + + if (!out) + out = "/tmp/egl_frame.ppm"; + f = fopen(out, "wb"); + if (!f) { + printf("PPM_OPEN_FAIL\n"); + return 1; + } + fprintf(f, "P6\n%d %d\n255\n", W, H); + for (i = 0; i < W * H; i++) + fwrite(px + i * 4, 1, 3, f); + fclose(f); + printf("EGL_FRAME_WRITTEN %s\n", out); + } + printf("EGL_FRAME_OK\n"); + return 0; +} + +int main(void) +{ + EGLDisplay dpy; + const char *ext; + + setvbuf(stdout, NULL, _IONBF, 0); + ext = eglQueryString(EGL_NO_DISPLAY, EGL_EXTENSIONS); + printf("EGL_CLIENT_EXT %s\n", ext ? ext : "(none)"); + dpy = get_display(); + if (dpy == EGL_NO_DISPLAY) { + printf("EGL_ALL_PLATFORMS_FAILED\n"); + return 1; + } + return render(dpy); +} diff --git a/tools/guest-diag/guest_pin_stream.py b/tools/guest-diag/guest_pin_stream.py index 97c68f2..175ed23 100644 --- a/tools/guest-diag/guest_pin_stream.py +++ b/tools/guest-diag/guest_pin_stream.py @@ -266,6 +266,103 @@ def _vmalloc_pages_for(kc, va): return 0 +def _nv_vmap_areas(kc): + """nv_vm_map_pages vmalloc areas: [(lo, hi, pages)] (5 s cache). + + Every CE push-buffer rm_mem is a vmap created by nv_vm_map_pages; the + demand pinner walks these areas to find the page holding a stalled + push VA by content, without needing any kprobe event or object chain. + """ + now = time.time() + if getattr(_nv_vmap_areas, "ts", 0) > now - 5.0: + return _nv_vmap_areas.areas + areas = [] + try: + for line in open("/proc/vmallocinfo"): + if "nv_vm_map_pages" not in line: + continue + p = line.split() + if len(p) < 2: + continue + try: + lo, hi = p[0].split("-") + lo = int(lo, 16) + hi = int(hi, 16) + except ValueError: + continue + pages = (hi - lo) >> 12 + if pages > 0: + areas.append((lo, hi, pages)) + except OSError: + pass + _nv_vmap_areas.areas = areas + _nv_vmap_areas.ts = now + return areas + + +# Learned page-offset per nv_vm_map_pages area for the demand content match +# (2026-10-16b): the first successful resolution records which page index the +# VA maps to, and later pushes in the same allocation retry that offset first +# instead of walking the whole +/-16 neighbourhood. +_area_off = {} + + +def resolve_va_by_content(kc, va, out): + """Demand pin: find the page carrying GPU VA 'va' by content match. + + The host reports the CE push-buffer VAs whose GPFIFO entries it cannot + read (dmesg 'stalled VA #N: 0x...'). The guest walks every + nv_vm_map_pages vmap area and tests the candidate leaf page at the VA's + in-window offset (plus a neighbourhood, for the per-mem header-page + offset ambiguity) against the host's method-stream prefix check. On a + match it emits the exact pair, so the host can translate and consume the + entry - no kprobe event, no object-chain discovery, no resolution + lottery. Only the requested VA is pinned, so a wrong candidate cannot + poison other VAs. + + 2026-10-16b: the offset search window was +/-1 page. Pushes of a client + whose pushbuffer vmap carries a larger per-allocation header (observed + with multi-MB CE staging: pushes at 0x120000f58+ never resolved and the + demand loop re-sent the same VA for minutes while the client crawled) + fall outside it. Widen to +/-16 pages and LEARN the winning offset per + vmap area, because consecutive pushes advance sequentially inside the + same allocation and share the offset. + """ + if not (PB_CE_LO <= va < PB_CE_HI): + return False + off = (va - PB_CE_LO) & 0xFFFFFF + k = off >> 12 + areas = _nv_vmap_areas(kc) + for (lo, _hi, pages) in areas: + learned = _area_off.get(lo) + if learned is not None: + adjs = [learned] + else: + adjs = [0, 1, -1] + adjs += [d for m in (2, 4, 8, 16) for d in (m, -m)] + for adj in adjs: + idx = k + adj + if idx < 0 or idx >= pages: + continue + cpu = lo + idx * 0x1000 + d = kc.read(cpu, 64) + if not d or not any(d[:32]): + continue + if not _push_prefix_ok(d): + continue + pa = kc.va_to_pa(cpu) + if pa is None: + continue + _area_off[lo] = adj + out.append("0x%x,0x%x,1\n" % (va & ~0xFFF, pa)) + dbg("demand-va: 0x%x -> pa=0x%x (area off=%d adj=%+d)" + % (va & ~0xFFF, pa, idx, adj)) + return True + dbg("demand-va: 0x%x not found by content (%d areas)" + % (va & ~0xFFF, len(areas))) + return False + + def resolve_mem_pins(kc, mem, out, default_pages=16, va_lo=0, va_hi=0, page_off=0): """uvm_rm_mem kptr -> GPU-VA/GPA runs (direct map or page-table walk). @@ -400,6 +497,121 @@ def mem_page_off(kc, mem): return 0 +def _ce_mem_pages(kc, mem): + """CE mem walk parameters: (gpus, cpu, off, npages) or None. + + Walks the same page set resolve_mem_pins used: the whole vmalloc area + (minus the header offset) when known, else size at +0x128, else the + default. The area governs: the live pushes are written sequentially and + may run past the size field, so a size-derived prefix stopped one page + short of the write frontier (r144: last page 0x120030000 emitted, the + live page 0x120031000 never in any pin set). all CE-window GPU VAs share + the CPU mapping. + """ + vas = [kc.read64(mem + 0x10 + i * 8) for i in range(16)] + cpu = next((v for v in vas if is_kern_ptr(v)), None) + gpus = [v for v in vas if v and PB_CE_LO <= v < PB_CE_HI] + if not cpu or not gpus: + return None + if _direct_map_pa(kc, cpu) is not None: + return None # never a vmap: not the push-buffer shape + off = mem_page_off(kc, mem) + area = _vmalloc_pages_for(kc, cpu) + size = _mem_size(kc, mem) + size_pages = max(1, min(size >> 12, 4098)) if size else 0 + area_pages = area - off if area and area > off else 0 + npages = min(max(size_pages, area_pages) or 64, 4098) + return (gpus, cpu, off, npages) + + +def _emit_ce_mem_delta(kc, st, out, fresh_map=False): + """Emit the CE mem's live {va,pa} pairs as merged runs. + + Two emission modes share one merge step: + + * full re-assert (KC_FULLREM_S cadence, and every fresh_map pass): every + page that translates is emitted, content-independent. This is the + fix for the host's session flush: at a client switch the host wipes + its whole sema map when the previous client's last channel closes + ("session flush (all channels closed)"), and a pin batch that raced + in for the new client is lost. The old content-gated delta never + re-emitted those unchanged pairs, so the new client's CE pushes were + permanently unreadable ("NO ENTRY" in sema-probe) and every GPFIFO + entry was skipped after the 8 s grace (r144/r145 second-client stall, + both mail and matmul). Re-asserting the full pair set every couple of + seconds heals any number of flush wipes; merged contiguous runs keep + the traffic to a few lines per pass. Zero-content pages are pinned + deliberately: the pin carries the translation, the host reads the + live content, and the client writes it later. + + * delta (between full passes): content-gated, emits pages whose PA is + new or changed, so a remap or a freshly-faulted page is picked up + without waiting for the next full pass. Never drops a previously + emitted pair (the host's last-store-wins pin sees the truth, and a + failed burst never clears the whole map). + """ + gpus, cpu, off, npages = st["gpus"], st["cpu"], st["page_off"], st["npages"] + last = st["last"] + nowm = time.monotonic() + full = nowm >= st.get("full_next", 0.0) + if full: + age = nowm - st.get("resolved_at", nowm) + st["full_next"] = nowm + (KC_FULLREM_S if age < KC_FULLREM_FAST_S + else KC_FULLREM_SLOW_S) + if fresh_map: + for k in range(npages): + kc._pa_cache.pop(cpu + (k + off) * 0x1000, None) + changed = [] + if full: + # Re-walk pages that did not translate on an earlier pass: the + # vmap's leaf sizes are huge (2 MB/1 GB), so most pages translate + # immediately; a straggler (the vmap populated lazily) is retried + # with its negative PA-cache entry popped, bounded per pass so the + # sweep stays cheap. + miss = st.setdefault("miss", set()) + retry = 0 + for k in range(npages): + page = cpu + (k + off) * 0x1000 + pa = kc.va_to_pa(page) + if pa is None: + if k in miss and retry < 512: + kc._pa_cache.pop(page, None) + retry += 1 + else: + miss.add(k) + continue + miss.discard(k) + for gv in gpus: + changed.append((gv + k * 0x1000, pa)) + else: + for k in range(npages): + d = kc.read(cpu + (k + off) * 0x1000, 32) + if not d or not any(d): + continue # zero/unreadable: nothing to pin yet + pa = kc.va_to_pa(cpu + (k + off) * 0x1000) + if pa is None: + continue + for gv in gpus: + if last.get(gv + k * 0x1000) != pa: + changed.append((gv + k * 0x1000, pa)) + changed.sort() + merged = [] + for va, pa in changed: + if merged and va == merged[-1][0] + merged[-1][2] * 0x1000 and \ + pa == merged[-1][1] + merged[-1][2] * 0x1000: + merged[-1][2] += 1 + else: + merged.append([va, pa, 1]) + for va, pa, n in merged: + out.append("0x%x,0x%x,%d\n" % (va, pa, n)) + for va, pa in changed: + last[va] = pa + st["emit_total"] = st.get("emit_total", 0) + len(merged) + if full: + st["full_total"] = st.get("full_total", 0) + 1 + return len(merged) + + def _emit_ring_pins(kc, gpu_va, entries_cpu, entries, out): """Pin the GPFIFO ring pages (and the pages after it, where the GPPUT buffer usually sits) from the ring's CPU mapping.""" @@ -522,6 +734,49 @@ def resolve_sem_pins(kc, sem, out): return resolve_mem_pins(kc, mem, out, default_pages=16) +def resolve_pool_va(kc, va, out, sems): + """Pool/waiter VA -> PA via tracker-sem -> page -> mem; NO content match. + + The word at POOL base + 4*idx is a slot in a UVM semaphore-pool page: a + sysmem uvm_rm_mem whose PA is only reachable through its kernel CPU + alias (the client pagemap has no present entry for it). Same kcore walk + the CE pushbuffer uses; the chain is shorter (sem+0 = page, page+0 = + mem). The page sits 16 MB above the CE push buffer base, outside + resolve_sem_pins' default_pages window, which is why the host saw + 'pending release starved' for it on every client's first seconds. + """ + page = va & ~0xFFF + if not (POOL_SEM_LO <= page < POOL_SEM_HI): + return False + for sem in list(sems): + if not is_kern_ptr(sem): + continue + try: + pg = kc.read64(sem) + if not is_kern_ptr(pg): + continue + mem = kc.read64(pg) + if not is_kern_ptr(mem): + continue + got = [] + resolve_mem_pins(kc, mem, got, default_pages=1, + va_lo=page, va_hi=page + 0x1000) + except Exception: + continue + for line in got: + try: + lva = int(line.split(',')[0], 16) + except Exception: + continue + if lva == page: + out.append(line) + dbg('pool-va: 0x%x -> %s (mem=0x%x)' + % (page, line.strip(), mem)) + return True + dbg('pool-va: 0x%x unresolved (%d known sems)' % (page, len(sems))) + return False + + def resolve_pb_pins(kc, pb, out): mem = kc.read64(pb + 8) if not mem: @@ -555,12 +810,35 @@ def resolve_pb_pins(kc, pb, out): # hold the ring-entry pushes is chosen (content anchor, _mem_content_ok). PB_CE_LO = 0x120000000 PB_CE_HI = 0x121000000 +# m165 (D-F1a): UVM semaphore-pool page range. The pool is a 4 KB sysmem +# uvm_rm_mem placed by the UVM arena right after the 16 MB CE push buffer; +# the client's tracker words (0x121000000 + 4*idx) live in its first page. +# The client pagemap has no present entry for it, so only the kcore +# tracker-sem -> page -> mem walk can translate it. +POOL_SEM_LO = 0x121000000 +POOL_SEM_HI = 0x122000000 +KC_POOL_REEMIT_S = 2.0 # re-assert the pool page pair this often too KC_CHAIN_BACK = 0x2000 KC_CHAIN_SPAN = 0xA000 KC_CHAIN_MAX = 400 RESOLVED_KC_PB = set() -KC_PB_STATE = {} # sem kptr -> resolved-mem fingerprint (freshness) +KC_PB_STATE = {} # sem kptr -> resolved-mem state (fingerprint + pins) KC_REFRESH_S = 2.0 # re-validate a resolved CE mem at most this often +KC_REEMIT_S = 2.0 # sweep the CE mem's live pages this often +KC_REMAP_S = 30.0 # full re-translation sweep (catches remaps) +# Full re-assert cadence: every pass at least this far apart re-emits the +# whole translatable pair set regardless of the delta, healing host-side +# session-flush wipes (see _emit_ce_mem_delta). Must stay well under the +# host's gsp_sema_ce_grace_ms (8 s) so a wiped-and-needed page is restored +# inside the hold window. The fast cadence applies to a mem's first +# KC_FULLREM_FAST_S seconds (the flush race lives at the client switch); +# after that the re-assert backs off to KC_FULLREM_SLOW_S - a fully +# fragmented pair set (descending PAs) emits one page per line, and 4 k +# lines every 2 s saturated a guest under multi-client load (r148: load 19, +# three scanners from an install race, clients wedged). +KC_FULLREM_S = float(os.environ.get("STREAM_FULLREM_S", "1.5")) +KC_FULLREM_FAST_S = float(os.environ.get("STREAM_FULLREM_FAST_S", "45")) +KC_FULLREM_SLOW_S = float(os.environ.get("STREAM_FULLREM_SLOW_S", "10.0")) def _chain_mems_at(kc, w): @@ -639,6 +917,16 @@ def _kc_pb_fresh(kc, sem): return False if not st.get("page_va"): return True + if not st.get("page_pa"): + # Unknown PA at store time: the manual page-table walk could not + # translate the vmap page (it may have been read before the kernel + # page tables were walkable). Treat the record as fresh instead of + # failing every check - a permanently-failing freshness test made + # the scanner drop the latch and re-resolve in a tight loop, and + # each failed re-resolution left the host with no pins at all + # (observed r131: "kc-chain stale: re-resolving" every second for + # five sems while smiley sat in cudaMalloc). + return True pa = kc.va_to_pa(st["page_va"]) if pa is None or pa != st.get("page_pa"): return False @@ -669,6 +957,42 @@ def resolve_pb_chain(kc, sem, out): return False words = struct.unpack_from("<%dQ" % (len(data) // 8), data, 0) cands = list(dict.fromkeys(w for w in words if is_kern_ptr(w))) + # Channel-anchored expansion (r134 live dump): the tracking sem is a + # member of a uvm_channel_t table; the channel records sit in the sem's + # own window and carry their CE ring GPU VA ([0x121010000, 0x121400000)) + # plus back-to-back 32-bit object handles. The pool/manager/pushbuffer + # pointers hang off those records' targets, which the bare sem-window + # scan never reaches - the second-client stall is exactly this: the only + # live uvmt event is the spinning sem, and no mem candidate is found + # around it. For each ring-VA word, scan the surrounding record for + # kptrs, take those targets, and try the chain shape on both. Additive + # and bounded; content scoring still decides, so a wrong candidate + # cannot shadow the live pages. + extra = [] + for i, w in enumerate(words): + if not w or not (0x121010000 <= w < 0x121400000): + continue + lo = max(0, i - 0x60) + hi = min(len(words), i + 0x60) + for c in words[lo:hi]: + if not is_kern_ptr(c): + continue + extra.append(c) + tgt = kc.read(c, 0x100) + if tgt: + tw = struct.unpack_from("<%dQ" % (len(tgt) // 8), + tgt, 0) + for w2 in tw: + if is_kern_ptr(w2): + extra.append(w2) + if len(extra) >= 512: + break + if extra: + head = [] + for w in extra: + if w not in head and w not in cands: + head.append(w) + cands = head + cands mems = [] for w in cands[:KC_CHAIN_MAX]: for mem in _chain_mems_at(kc, w): @@ -697,18 +1021,24 @@ def resolve_pb_chain(kc, sem, out): if mem in RESOLVED_KC_PB: return True RESOLVED_KC_PB.add(mem) - off = mem_page_off(kc, mem) - n = resolve_mem_pins(kc, mem, out, default_pages=64, - va_lo=PB_CE_LO, va_hi=PB_CE_HI, - page_off=off) - st = {"mem": mem, "gpu": gpu, "cpu": cpu, - "page_va": cpu, "page_pa": kc.va_to_pa(cpu) or 0, "page_off": 0} + params = _ce_mem_pages(kc, mem) + if not params: + return found + gpus, cpu, off, npages = params + st = {"mem": mem, "gpu": gpu, "cpu": cpu, "gpus": gpus, + "page_va": cpu, "page_pa": kc.va_to_pa(cpu) or 0, + "page_off": off, "npages": npages, "last": {}, + "resolved_at": time.monotonic()} + # Fresh translation for the first sweep: the Kcore PA cache may hold + # pairs from before this mem's vmap area was (re)populated. + kc._pa_cache.clear() + n = _emit_ce_mem_delta(kc, st, out) KC_PB_STATE[sem] = st try: with open('/tmp/guest_pin_stream.log', 'a') as f: f.write('%s kc-chain CE push-buffer mem=0x%x sem=0x%x ' - 'pin-runs=%d off=%d size=%d score=%d\n' - % (time.strftime('%T'), mem, sem, n, off, + 'pin-runs=%d off=%d npages=%d size=%d score=%d\n' + % (time.strftime('%T'), mem, sem, n, off, npages, _mem_size(kc, mem), _score)) except OSError: pass @@ -975,6 +1305,75 @@ def new_sems(seen): return found +# --------------------------------------------------------------- scrubber +# PMA scrubber channel discovery (third-client scrub freeze). The guest's +# PMA scrubber is an internal CeUtils CE channel (4096 GPFIFO entries, +# pbGpuVA 0x120000000, ring at pbGpuVA+0x64000, finish payload at +# pbGpuVA+0x6c004). It lives in a kernel vmap area, not in any user pagemap, +# and its GPU VAs collide numerically with the client CE push-buffer window +# in the host's global translation map - so the host cannot resolve it and +# must be told the finish/host-semaphore word PAs directly. Resolve them +# here (vmallocinfo + manual kernel page-table walk) and emit one +# #scrub <fin_pa>,<hsem_pa> +# line; pin_gr_stream.py registers them through sema_scrubbuf_add. +SCRUB_AREA_SIZES = (0x6e000, 0x6d000, 0x6c000, 0x6b000, 0x6f000, + 0x70000, 0x71000, 0x72000) +_scrub_sent = {} + + +def _scrub_ring_ok(kc, ring_pa): + """A 906F GPFIFO ring page: all-zero (fresh) or entries whose GET names a + push inside the channel's own 0x120000000 window (encoded as + push_va - 0x100000000, 4-byte aligned).""" + buf = kc.read(kc.page_offset_base + ring_pa, 0x1000) + if not buf: + return False + for off in range(0, 0x100, 8): + e0, _e1 = struct.unpack_from('<II', buf, off) + if e0 == 0: + continue + if not (0x20000000 <= e0 < 0x20080000) or (e0 & 3): + return False + return True + + +def scrub_scan(kc, lines): + """Emit the scrubber buffer words' PAs once (re-emitted if they move).""" + try: + f = open('/proc/vmallocinfo') + except OSError: + return + with f: + for line in f: + if 'nvidia' not in line or 'ioremap' in line or 'vmap' not in line: + continue + parts = line.split() + try: + a, b = parts[0].split('-') + lo, hi = int(a, 16), int(b, 16) + except (ValueError, IndexError): + continue + if (hi - lo) not in SCRUB_AREA_SIZES: + continue + fin_pa = kc.va_to_pa(lo + 0x6c000) + ring_pa = kc.va_to_pa(lo + 0x64000) + if fin_pa is None or ring_pa is None: + continue + if not _scrub_ring_ok(kc, ring_pa): + dbg('scrub candidate 0x%x rejected (ring shape)' % lo) + continue + fin = fin_pa + 4 + hsem = fin_pa + key = (fin, hsem) + if _scrub_sent.get('key') == key: + return + _scrub_sent['key'] = key + lines.append('#scrub 0x%x,0x%x\n' % (fin, hsem)) + dbg('scrub buffer emitted lo=0x%x fin=0x%x hsem=0x%x' + % (lo, fin, hsem)) + return + + def nvidia_fds(pid): """CUDA device fds this pid holds: 'uvm', 'gpu', or None. @@ -999,6 +1398,34 @@ def nvidia_fds(pid): return kind +def pid_starttime(pid): + """Starttime (clock ticks, field 22) of /proc/<pid>, or None. + + Part of the client identity (pid, starttime): pids are recycled, so a + per-client reset keyed on the pid alone can mis-fire on a recycled pid. + Parsed after the last ')' so comm names with spaces/parens are safe.""" + try: + with open('/proc/%d/stat' % pid) as f: + st = f.read() + rp = st.rfind(')') + if rp < 0: + return None + toks = st[rp + 2:].split() + if len(toks) < 20: # starttime = field 22 overall = index 19 + return None + return int(toks[19]) + except (OSError, ValueError): + return None + + +def pid_line(tag, pid, start): + """'#<tag> <pid> [starttime]' - the optional starttime keeps the line + backward compatible with older hosts (they parse the pid only).""" + if start is None: + return "#%s %d\n" % (tag, pid) + return "#%s %d %d\n" % (tag, pid, start) + + def cdt_pid(): prefixes = tuple(p for p in os.environ.get("CLIENT_PREFIXES", "cdt,va,victory").split(",") if p) @@ -1117,6 +1544,18 @@ def scan_vmas(fd, pid, out): parts = line.split() if len(parts) < 5: continue + # Skip mappings with no read permission: a CUDA device buffer + # is a PROT_NONE reservation (e.g. `---p 00000000 00:00 0`), + # potentially hundreds of MB. Its pages are never present in + # the pagemap, so scanning it produces no pairs - but it + # consumed the VMA_MAX_PAGES budget and the scan never reached + # the client's present staging buffers that sit ABOVE the + # device reservation (observed: a client whose device range + # started below its staging buffer got no pins at all and its + # H2D copies failed with rc=-2 forever). + perms = parts[1] if len(parts) >= 2 else '' + if 'r' not in perms: + continue # Anonymous mappings (no path field) are the client's UVM sysmem # buffers when the driver maps them without a /dev/zero backing: # the second client's cudaMalloc'd data buffers show up exactly @@ -1206,6 +1645,8 @@ def main(): resolved_kc_pb = set() retry_at = {} kc_refresh_at = {} + kc_reemit_at = {} + kc_remap_at = {} tracer = TraceReader() kc = Kcore() kc_ok = kc.load() @@ -1216,13 +1657,34 @@ def main(): except OSError: pass last_pid = None + last_start = None last_log = 0.0 + scrub_next = 0.0 + demand_next = 0.0 + # m165 (D-F1a): pool page VAs learned from demand, re-asserted every + # KC_POOL_REEMIT_S. Boot-stable (UVM frees pool pages only at GPU + # deinit), so the set deliberately persists across client switches. + pool_pages = {} + pool_reemit_next = 0.0 + _want_state = {"mtime": None, "done": set()} loop = 0 while True: pid = cdt_pid() loop += 1 if pid is None: + if last_pid is not None and sock is not None: + # In-band client-death signal: the host tools purge the dead + # client's per-client state on this line, so the NEXT client + # works without a VM cycle (the kill cascade). Best-effort - + # if the socket is gone the host still has its own probe. + try: + sock.sendall(pid_line("nopid", last_pid, + last_start).encode()) + except OSError: + sock.close() + sock = None last_pid = None + last_start = None seen = {} seen_sems = set() seen_pbs = set() @@ -1231,6 +1693,7 @@ def main(): continue if pid != last_pid: last_pid = pid + last_start = pid_starttime(pid) seen = {} seen_sems = set() seen_pbs = set() @@ -1242,6 +1705,8 @@ def main(): resolved_kc_pb = set() retry_at = {} kc_refresh_at = {} + kc_reemit_at = {} + kc_remap_at = {} # Client switch: the old client's channel/put pointers are freed # memory. Clearing the #put map stops stale reads from being # emitted as if they were the new client's counters. @@ -1259,7 +1724,7 @@ def main(): dbg("switch seed: +%d event kptr(s)" % _seeded) if sock is not None: try: - sock.sendall(("#pid %d\n" % pid).encode()) + sock.sendall(pid_line("pid", pid, last_start).encode()) except OSError: sock.close() sock = None @@ -1293,6 +1758,123 @@ def main(): new = [(va, pa) for va, pa in cur.items() if seen.get(va) != pa] pending_cur = cur lines = [] + # Demand pins (host -> this scanner): the host appends the CE + # push-buffer VAs of GPFIFO entries it cannot read to + # /tmp/guest_want_vas.txt (written via the guest agent by + # pin_gr_stream); resolve each by content against the + # nv_vm_map_pages areas and emit the pair. This is the + # event-independent hole filler for the stall class where the + # one-shot #sem kprobe event was lost and the object-chain + # resolver never learns the client. + if kc_ok and time.time() >= demand_next: + demand_next = time.time() + 1.0 + try: + st_d = os.stat("/tmp/guest_want_vas.txt") + if st_d.st_mtime != _want_state.get("mtime"): + with open("/tmp/guest_want_vas.txt") as fw: + # r582: tolerate malformed historical lines. The + # r581 boot's scanner died here: a transport-escaped + # append merged VAs into one token containing + # literal backslash-n ('0x20400f000\\n0x121000000\\n + # ...') and int(t, 16) raised ValueError out of + # main(), stopping ALL new pins (the observed + # pool-pin latency). Extract every 0x token instead; + # a bad token is skipped, never fatal. + wvas = [] + for t in re.findall(r"0x[0-9a-fA-F]+", + fw.read()): + try: + wvas.append(int(t, 16)) + except ValueError: + pass + _want_state["mtime"] = st_d.st_mtime + for wv in wvas[-64:]: + if wv in _want_state["done"]: + continue + got = [] + if POOL_SEM_LO <= (wv & ~0xFFF) < POOL_SEM_HI: + # m165 (D-F1a): pool/waiter VAs have no method + # stream content and the CE offset math does not + # apply; route them to the structural tracker-sem + # walk instead. m168: the tracker sems are + # long-lived objects (allocated once) whose page + # pointer follows the pool across clients, so try + # every sem we have ever resolved a chain for - + # seen_sems alone can miss the live generation. + _psems = set(seen_sems) | set(KC_PB_STATE.keys()) \ + | set(resolved_kc_pb) + try: + ok = resolve_pool_va(kc, wv, got, + _psems) + except Exception: + ok = False + got = [] + dbg('demand-pool-va: 0x%x %s' + % (wv, 'ok' if ok else 'pending')) + if ok and got: + pool_pages.setdefault(wv & ~0xFFF, True) + else: + try: + ok = resolve_va_by_content(kc, wv, got) + except Exception: + ok = False + got = [] + # r586: a demanded VA outside the CE + # pushbuffer window (the GRFILL staging pages, + # high client VAs) is a plain client mapping + # page: resolve it through the client's own + # pagemap. Without this the demand for those + # VAs was a silent no-op (resolve_va_by_content + # rejects anything outside PB_CE_LO..PB_CE_HI) + # and the fills kept dropping them. + if not ok and not got and pid: + try: + pfd = os.open('/proc/%d/pagemap' % pid, + os.O_RDONLY) + try: + sub = {} + vpg = wv & ~0xFFF + scan_range(pfd, vpg, vpg + 0x1000, + sub) + finally: + os.close(pfd) + if sub: + for v, p in sub.items(): + got.append('0x%x,0x%x,1\n' + % (v, p)) + ok = True + except OSError: + pass + if ok and got: + lines.extend(got) + _want_state["done"].add(wv) + except OSError: + pass + # m165 (D-F1a): re-assert the pool page pair every KC_POOL_REEMIT_S. + # The pool page is boot-stable (UVM frees it only at GPU deinit) but + # the host pin table is wiped at every session flush, so a one-shot + # pin leaves later clients starving until the slow structural walks + # answer; the periodic re-assert heals the wipe within seconds (the + # same argument as the CE mem sweep above). + if kc_ok and pool_pages and time.time() >= pool_reemit_next: + pool_reemit_next = time.time() + KC_POOL_REEMIT_S + for pv in list(pool_pages): + got = [] + try: + ok = resolve_pool_va(kc, pv, got, seen_sems) + except Exception: + ok = False + got = [] + if ok and got: + lines.extend(got) + # PMA scrubber buffer discovery (throttled): resolve the finish/host- + # semaphore word PAs and stream them for host registration. + if kc_ok and time.time() >= scrub_next: + scrub_next = time.time() + 5.0 + try: + scrub_scan(kc, lines) + except Exception as e: + dbg('scrub scan failed: %s' % e) for ctl in tracer.drain(): parts = ctl.split() if ctl.startswith("#sem") and len(parts) >= 2: @@ -1333,12 +1915,68 @@ def main(): except Exception: _fresh = True if not _fresh: - _st = KC_PB_STATE.pop(sem, None) + # Re-resolve, but never lose the latch on a failed + # attempt: a dropped latch means the host keeps no + # pins for this channel at all, which is worse than + # a possibly-stale pin (the host content-checks the + # first page and now rejects all-FF continuation + # chunks, so stale pages are bounded). On failure + # keep the old state and defer the next attempt. + _st = KC_PB_STATE.get(sem) if _st and _st.get("mem"): RESOLVED_KC_PB.discard(_st["mem"]) - resolved_kc_pb.discard(sem) - retry_at.pop(("pk", sem), None) - dbg("kc-chain stale: re-resolving sem=0x%x" % sem) + got = [] + try: + _found2 = resolve_pb_chain(kc, sem, got) + except Exception: + _found2 = False + got = [] + if _found2 and got: + _new = KC_PB_STATE.get(sem) + _newmem = _new.get("mem") if _new else None + dbg("kc-chain stale: re-resolved sem=0x%x " + "new-mem=%s runs=%d" + % (sem, "yes" if _newmem != + (_st.get("mem") if _st else None) + else "no", len(got))) + lines.extend(got) + added += len(got) + resolved_kc_pb.add(sem) + else: + if _st and _st.get("mem"): + RESOLVED_KC_PB.add(_st["mem"]) + if sem not in KC_PB_STATE: + KC_PB_STATE[sem] = _st + kc_refresh_at[sem] = now_r + 5.0 + dbg("kc-chain stale: keep old latch sem=0x%x" + % sem) + else: + # CE mem delta sweep (second-client pins): the one-shot + # emission pinned the mem once; pages the client writes + # later (or whose vmap pairing moved) were never + # re-emitted, so the host kept zeros for the other CE + # ranks' pushes. Sweep every KC_REEMIT_S, emitting only + # new/changed nonzero pairs; burst the Kcore PA cache + # every KC_REMAP_S so a remap is not shadowed. + _st = KC_PB_STATE.get(sem) + if _st is not None and \ + now_r >= kc_reemit_at.get(sem, 0): + _fresh_map = now_r >= kc_remap_at.get(sem, 0) + if _fresh_map: + kc_remap_at[sem] = now_r + KC_REMAP_S + kc_reemit_at[sem] = now_r + KC_REEMIT_S + got = [] + try: + _n2 = _emit_ce_mem_delta(kc, _st, got, + fresh_map=_fresh_map) + except Exception: + _n2 = 0 + got = [] + if got: + lines.extend(got) + added += len(got) + dbg("ce-delta sem=0x%x runs=%d mem=0x%x" + % (sem, _n2, _st.get("mem", 0))) # CE push-buffer chain (second-client stall fix): resolve # pool -> manager -> pushbuffer -> rm_mem directly, so the # push-buffer pages never depend on the racy one-shot uvbp @@ -1480,7 +2118,8 @@ def main(): sock.connect((CID, PORT)) sock.settimeout(None) if last_pid is not None: - sock.sendall(("#pid %d\n" % last_pid).encode()) + sock.sendall(pid_line("pid", last_pid, + last_start).encode()) except OSError: if sock: sock.close() diff --git a/tools/guest-diag/pin_gr_stream.py b/tools/guest-diag/pin_gr_stream.py index b864ab5..ff7e299 100644 --- a/tools/guest-diag/pin_gr_stream.py +++ b/tools/guest-diag/pin_gr_stream.py @@ -16,6 +16,7 @@ import re import socket import subprocess import sys +import threading import time HERE = os.path.dirname(os.path.abspath(__file__)) @@ -31,12 +32,44 @@ LO = os.environ.get("PINWIN_LO", "0x200000000") HI = os.environ.get("PINWIN_HI", "0x206000000") PAMAX = os.environ.get("PINWIN_PA_MAX", "0x180000000") SEMSTREAM = os.environ.get("SEM_STREAM_FILE", "/tmp/opencode/sem_stream.txt") +# Liveness file checked by pin_supervisor.sh (mtime must stay < 75 s; the +# supervisor grants a 120 s startup grace on top). +HEARTBEAT_FILE = os.environ.get("PIN_GR_HEARTBEAT", "/tmp/opencode/pin_gr_heartbeat") +# Inter-tool re-assert request (written by semfix_uvmt.py's stall watchdog): +# when its mtime moves, re-assert every pair this tool has ever pinned. +REASSERT_REQ = os.environ.get("PIN_REASSERT_FILE", + "/tmp/opencode/pin_reassert_request") +# Stall watchdog: a known client with no pin/data activity for this long gets +# an automatic recovery (full re-assert + demand of the fixed counter/pool +# VAs). Pre-capture (no pin yet for this client) uses the tighter limit. +STALL_PRE_S = float(os.environ.get("PIN_STALL_PRE", "60")) +STALL_POST_S = float(os.environ.get("PIN_STALL_POST", "120")) +STALL_RETRY_S = float(os.environ.get("PIN_STALL_RETRY", "30")) +STALL_ESCALATE_S = float(os.environ.get("PIN_STALL_ESCALATE", "180")) def log(msg): print("[%s] %s" % (time.strftime("%H:%M:%S"), msg), flush=True) +def _heartbeat(): + try: + with open(HEARTBEAT_FILE, "w") as f: + f.write("%.3f\n" % time.time()) + except OSError: + pass + + +def _hb_loop(): + """Refresh the heartbeat on a fixed cadence, independent of data-path + activity: recv can be legally idle for minutes between client phases and + the supervisor used to read that silence as a hung tool (r198 restarts at + 34 s staleness while injection bursts were running).""" + while True: + _heartbeat() + time.sleep(5) + + def _mirror(line): """Append guest control lines (#pid/#sem) to the shared host file so the single-QMP watcher can consume them without guest-agent round trips.""" @@ -56,18 +89,68 @@ def vexec(cmd, timeout=40): return "" +def _scanner_alive(): + """True when the in-guest scanner process still exists. + + install_and_start() pkill's and restarts the scanner, which resets every + in-guest latch; probing first avoids killing a healthy scanner that is + merely silent because no client exists yet (r198 pre-client churn).""" + out = vexec("sudo sh -c 'kill -0 $(cat /tmp/guest_pin_stream.pid) " + "2>/dev/null && echo ALIVE'") + return "ALIVE" in (out or "") + + +def _demand_write_cmd(vas): + """Build a transport-proof append command for the guest demand file. + + r582: the old `printf '%s' '<lines with real LF>'` form was mangled by + the guest-agent transport under load - the newlines arrived as LITERAL + backslash-n text and merged VAs onto one line, and the guest scanner's + parser then crashed on `int('0x20400f000\\\\n0x121000000...', 16)` + (ValueError in main; r581 boot scanner log). The scanner's death stopped + ALL new pins while the host only re-asserted known pairs - the observed + pool-pin-latency (0x121000000 never arrived while other injections + flowed). One `echo <va> >> file` per VA carries no newlines and no + backslashes at all, so no transport layer can escape them.""" + return "; ".join("echo 0x%x >> /tmp/guest_want_vas.txt" % va + for va in vas) + + +def _demand_send(vas): + """Append VA hints to the guest scanner's demand file (best-effort). + + Same channel pin_gr_stream._demand_loop uses for host-dmesg stalls; the + stall watchdog uses it to ask the scanner to re-resolve the fixed + pre-capture VAs (counter page / UVM pool) by content/pagemap.""" + cmd = _demand_write_cmd(vas) + if not cmd: + return + try: + vexec(cmd, 30) + except Exception: + pass + + def _start_cmd(): with open(GUEST_SCRIPT) as f: b64 = subprocess.run(["base64", "-w0"], stdin=f, - capture_output=True, text=True).stdout + capture_output=True, text=True, + timeout=10).stdout prefixes = os.environ.get("CLIENT_PREFIXES", "cdt,va,victory") vma_every = os.environ.get("STREAM_VMA_EVERY", "1") - return ("echo '%s' | base64 -d > /tmp/guest_pin_stream.py; " - "if [ -f /tmp/guest_pin_stream.pid ]; then " - "p=$(cat /tmp/guest_pin_stream.pid 2>/dev/null); " - "[ -n \"$p\" ] && grep -q guest_pin_stream " - "/proc/$p/cmdline 2>/dev/null && sudo kill $p 2>/dev/null; " - "fi; " + # Kill EVERY existing scanner, not just the pid-file one: a retry after + # a partially-failed first attempt (pid file not written yet) used to + # leave the first scanner alive, and multiple scanners each re-assert + # the whole CE pair set - r148 ran three and drove the guest to load 19 + # (clients wedged in userspace spins). The pattern is anchored + # (^python3) and assembled at runtime via printf: an unanchored pattern + # or a literal one would also match this command's own shell (its + # cmdline contains the script path), and pkill would kill the shell + # before it could start the new scanner - every start attempt then + # failed silently (r149). + return ("P=$(printf '^python3 /tmp/guest_pin_%%s.py' stream); " + "sudo pkill -f \"$P\" 2>/dev/null; sleep 1; " + "echo '%s' | base64 -d > /tmp/guest_pin_stream.py; " "sudo rm -f /tmp/guest_pin_stream.pid; " "sudo setsid nohup env CLIENT_PREFIXES='%s' " "STREAM_VMA_EVERY='%s' " @@ -82,6 +165,7 @@ def install_and_start(): # or a stale qga client); retry a few times instead of giving up, because # a boot without the scanner streams no pins and the client then hangs. for attempt in range(5): + _heartbeat() out = vexec(_start_cmd(), 60) if any(ch.isdigit() for ch in out): return True @@ -93,78 +177,297 @@ def install_and_start(): def inject(va, pa, ln): cmd = "printf '0x%x,0x%x,0x%x,1' > %s/sema_map_add" % (va, pa, ln, PARAM) for _ in range(60): - r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], - capture_output=True, text=True) + try: + r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], + capture_output=True, text=True, timeout=10) + except subprocess.TimeoutExpired: + time.sleep(0.05) + continue if r.returncode == 0: return True time.sleep(0.05) return False +def scrub_register(fin_pa, hsem_pa): + """Register the PMA scrubber buffer words (guest-resolved PAs) with the + service so it can author the scrub progress (third-client scrub freeze).""" + cmd = ("printf '0x%x,0x%x' > %s/sema_scrubbuf_add" + % (fin_pa, hsem_pa, PARAM)) + for _ in range(60): + try: + r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], + capture_output=True, text=True, timeout=10) + except subprocess.TimeoutExpired: + time.sleep(0.05) + continue + if r.returncode == 0: + log("scrub registered fin=0x%x hsem=0x%x" % (fin_pa, hsem_pa)) + return True + time.sleep(0.05) + log("scrub register FAILED fin=0x%x hsem=0x%x" % (fin_pa, hsem_pa)) + return False + + +def _demand_loop(): + """Push the host's unreadable CE push VAs to the guest scanner. + + The service logs each distinct stalled VA once per boot ('stalled VA + #N: 0x...'). Each sighting is tracked and re-sent every 20 s for five + minutes: the guest resolves the VA by content against the + nv_vm_map_pages areas (event-independent, no object chain), so a page + that was not yet written when first asked is picked up on a retry. + This closes the stall class where the one-shot #sem kprobe event was + lost and the chain resolver never learns the client. + """ + track = {} # va -> [first_seen, last_sent] + while True: + time.sleep(3) + try: + r = subprocess.run(["sudo", "-n", "dmesg"], + capture_output=True, text=True, timeout=20) + except Exception: + continue + now = time.time() + for m in re.finditer(r"(?:stalled VA #\d+[:=]\s*|pending release starved va=)(0x[0-9a-f]+)", + (r.stdout or "")[-200000:]): + va = int(m.group(1), 16) + track.setdefault(va, [now, 0.0]) + for va in list(track): + first, last = track[va] + if now - first > 300.0: + del track[va] + todo = [va for va, (f, l) in track.items() if now - l > 20.0] + if not todo: + continue + cmd = _demand_write_cmd(todo) + try: + vexec(cmd, 30) + except Exception: + continue + for va in todo: + track[va][1] = now + log("demand: sent %d stalled VA(s) to guest" % len(todo)) + + +_DEMAND_STARTED = False +_HB_STARTED = False +_REASSERT_REQ_SEEN = 0.0 + + def main(): + global _DEMAND_STARTED, _HB_STARTED + _heartbeat() + if not _HB_STARTED: + _HB_STARTED = True + try: + threading.Thread(target=_hb_loop, daemon=True).start() + except Exception as e: + log("WARN: heartbeat thread failed to start: %s" % e) if not install_and_start(): log("WARN: could not start guest scanner (will retry via vsock wait)") + if not _DEMAND_STARTED: + _DEMAND_STARTED = True + try: + threading.Thread(target=_demand_loop, daemon=True).start() + except Exception as e: + log("WARN: demand thread failed to start: %s" % e) srv = socket.socket(socket.AF_VSOCK, socket.SOCK_STREAM) srv.setsockopt(socket.SOL_SOCKET, socket.SO_REUSEADDR, 1) srv.bind((socket.VMADDR_CID_ANY, PORT)) + srv.settimeout(10) srv.listen(4) log("listening on vsock port %d (window %s..%s)" % (PORT, LO, HI)) seen = {} pin_total = 0 - buf = b"" + last_conn = time.time() + last_reinstall = 0.0 + # Stall-watchdog state for the connected client. client_pid is the + # scanner-reported pid (None until a client appears); last_activity marks + # the last pin/data/control event; pins_since_client counts pages pinned + # for the current client (0 = pre-capture). + client_pid = None + client_start = None + last_activity = time.time() + pins_since_client = 0 + stall_since = None + recover_next = 0.0 + escalated = False + no_conn_warn_next = 0.0 + + def _activity(): + nonlocal last_activity, stall_since, escalated + last_activity = time.time() + stall_since = None + escalated = False + while True: try: conn, _ = srv.accept() + except socket.timeout: + now = time.time() + # No guest scanner connection for 30 s: it died or never started + # (a boot without it streams no pins and every client then hangs). + # Probe the scanner first instead of reinstalling blindly: + # install_and_start() pkill's a live scanner and resets all + # in-guest latches, and pre-client silence is normal (r198 churn). + if now - last_conn > 30.0 and now - last_reinstall >= 60.0: + last_reinstall = now + if _scanner_alive(): + log("no guest connection for %.0fs but scanner alive; " + "not reinstalling" % (now - last_conn)) + if now - last_conn > 180.0 and now >= no_conn_warn_next: + no_conn_warn_next = now + 60.0 + log("STALL WATCHDOG: no guest scanner connection for " + "%.0fs (scanner process alive) - pins/demand may " + "be stranded" % (now - last_conn)) + continue + log("reinstall attempt: no guest connection for %.0fs" + % (now - last_conn)) + try: + install_and_start() + except Exception as e: + log("reinstall failed: %s" % e) + continue except OSError as e: log("accept failed: %s" % e) time.sleep(1) continue + last_conn = time.time() + _heartbeat() conn.settimeout(5) log("guest connected") buf = b"" runs = [] - while True: - try: - chunk = conn.recv(1 << 20) - except socket.timeout: + try: + while True: + try: + chunk = conn.recv(1 << 20) + except socket.timeout: + if runs: + novel = _novel_pages(runs, seen) + pages = _flush(runs, seen) + pin_total += pages + pins_since_client += pages + if novel: + _activity() + runs = [] + _maybe_reassert_request(seen) + now = time.time() + # Stall watchdog: a known client with no pin/data + # activity gets a bounded, idempotent recovery (re-assert + # everything + demand the fixed counter/pool VAs), then a + # loud escalation if the stall persists. This is the + # "first cudaMalloc / first GR capture must not stall + # silently" guard. recover_next rate-limits attempts (the + # loop wakes every 5 s on the recv timeout). + if client_pid is not None: + limit = STALL_PRE_S if pins_since_client == 0 \ + else STALL_POST_S + idle = now - last_activity + if idle >= limit and now >= recover_next: + recover_next = now + STALL_RETRY_S + if stall_since is None: + stall_since = now + log("STALL WATCHDOG: client pid %d - no pin " + "activity for %.0fs (pre-capture=%s) - " + "re-asserting all known pairs + demanding " + "counter/pool VAs" + % (client_pid, idle, pins_since_client == 0)) + if seen: + _flush([], seen, full=True) + _demand_send([0x20400f000, 0x121000000]) + if not escalated and stall_since is not None \ + and now - stall_since >= STALL_ESCALATE_S: + escalated = True + log("STALL WATCHDOG: recovery FAILED for pid %d " + "after %.0fs - stall persists, needs manual " + "attention" % (client_pid, now - stall_since)) + continue + except OSError: + break + if not chunk: + break + _heartbeat() + buf += chunk + while b"\n" in buf: + line, buf = buf.split(b"\n", 1) + text = line.decode(errors="replace").strip() + if not text: + continue + if text.startswith("#pid"): + m = re.match(r"#pid\s+(\d+)(?:\s+(\d+))?", text) + seen = {} + client_pid = int(m.group(1)) if m else None + client_start = (int(m.group(2)) if m and m.group(2) + else None) + pins_since_client = 0 + _activity() + log("guest reports cdt pid %s%s" + % (client_pid if client_pid is not None else "?", + (" start=%s" % client_start + if client_start is not None else ""))) + _mirror(text) + continue + if text.startswith("#nopid"): + # In-band client death: the guest scanner tells us the + # client is gone (with its identity). Drop the per-pair + # dedupe map so the next client's pairs are injected + # unconditionally; semfix_uvmt consumes the mirrored + # line and purges its per-client state too. + client_pid = None + client_start = None + seen = {} + pins_since_client = 0 + _activity() + log("guest reports client gone (%s)" % text) + _mirror(text) + continue + if text.startswith("#sem"): + _mirror(text) + _activity() + continue + if text.startswith("#pb"): + _mirror(text) + _activity() + continue + if text.startswith("#put"): + log(text) + _mirror(text) + _activity() + continue + if text.startswith("#scrub"): + m = re.search(r"0x([0-9a-fA-F]+),0x([0-9a-fA-F]+)", + text) + _mirror(text) + if m: + scrub_register(int(m.group(1), 16), + int(m.group(2), 16)) + continue + m = re.match(r"^(0x[0-9a-f]+),(0x[0-9a-f]+)(?:,(\d+))?$", + text) + if m: + runs.append((int(m.group(1), 16), + int(m.group(2), 16), + int(m.group(3)) if m.group(3) else 1)) if runs: - _flush(runs, seen) + novel = _novel_pages(runs, seen) + pages = _flush(runs, seen) + pin_total += pages + pins_since_client += pages + if novel: + _activity() runs = [] - continue - except OSError: - break - if not chunk: - break - buf += chunk - while b"\n" in buf: - line, buf = buf.split(b"\n", 1) - text = line.decode(errors="replace").strip() - if not text: - continue - if text.startswith("#pid"): - m = re.search(r"(\d+)", text) - seen = {} - log("guest reports cdt pid %s" % (m.group(1) if m else "?")) - _mirror(text) - continue - if text.startswith("#sem"): - _mirror(text) - continue - if text.startswith("#pb"): - _mirror(text) - continue - if text.startswith("#put"): - log(text) - _mirror(text) - continue - m = re.match(r"^(0x[0-9a-f]+),(0x[0-9a-f]+)(?:,(\d+))?$", text) - if m: - runs.append((int(m.group(1), 16), int(m.group(2), 16), - int(m.group(3)) if m.group(3) else 1)) - if runs: - pin_total += _flush(runs, seen) - runs = [] + # Poll the semfix re-assert request on every recv cycle, not + # only on recv-timeouts: sustained scanner traffic starves the + # timeout branch, so the request fired post-mortem in r203 + # (119 s late, `seen` already cleared by #nopid). + _maybe_reassert_request(seen) + except Exception as e: + # A transient error on one connection must not kill the process: + # log it, drop the connection and go back to accept(). + log("connection handler error: %s (continuing)" % e) log("guest disconnected (pins this session: %d)" % pin_total) try: conn.close() @@ -172,19 +475,82 @@ def main(): pass -def _flush(runs, seen): +def _maybe_reassert_request(seen): + """Honour a full-re-assert request written by semfix_uvmt's watchdog.""" + global _REASSERT_REQ_SEEN + try: + m = os.stat(REASSERT_REQ).st_mtime + except OSError: + return + if m > _REASSERT_REQ_SEEN: + _REASSERT_REQ_SEEN = m + if seen: + log("semfix requested a full pin re-assert - re-asserting %d " + "pair(s)" % len(seen)) + _flush([], seen, full=True) + + +def _coalesce_seen(seen): + """Coalesce the seen {va: pa} per-page map into (va, pa, npages) runs.""" + items = sorted(seen.items()) + runs = [] + i = 0 + while i < len(items): + va, pa = items[i] + n = 1 + while (i + n < len(items) + and items[i + n][0] == va + n * 0x1000 + and items[i + n][1] == pa + n * 0x1000): + n += 1 + runs.append((va, pa, n)) + i += n + return runs + + +def _novel_pages(runs, seen): + """Count pages in runs whose translation is new or changed. + + Progress accounting for the stall watchdog: the guest re-streams the CE + window and UVM pool every couple of seconds even when nothing moves, so a + pure repeat of an already-known pair must not count as pin activity + (r203: those repeats masked the watchdog through every hang).""" + return sum(1 for va, pa, n in runs for i in range(n) + if seen.get(va + i * 0x1000) != pa + i * 0x1000) + + +def _flush(runs, seen, full=False): """Inject run lines (one batched sudo pass), return number of pages pinned. The old implementation spawned one sudo process per run; with ~400 runs that cost ~4 s, which landed the GR pins after the client's context-init deadline. The batch writes all runs in a single shell, retrying only the ones the kernel queue rejected (ENOSPC) individually. + + full=True re-asserts every pair ever pinned (coalesced), not just this + batch's runs: the kernel session flush wipes the pin table when a client's + channels close, and the stall watchdog / semfix re-assert request use this + to heal it without touching the guest. Writes are idempotent. """ - todo = [] + todo = {} for va, pa, n in runs: - if seen.get(va) == pa and n == 1: + # The CE window (0x120000000..0x121000000) and the UVM semaphore + # pool page range (0x121000000..0x122000000) are re-asserted by the + # guest every couple of seconds on purpose: the host's session flush + # wipes the whole pin table when the previous client's last channel + # closes, and a pin batch that raced in for the new client must be + # allowed to land again even when the pair is unchanged. Skipping + # repeats there stranded the second client's CE pushes (the r144/ + # r145 stall: sema-probe "NO ENTRY" for the live push pages) and the + # pool/tracker page (pending release starved, every client's first + # seconds). + ce_win = 0x120000000 <= va < 0x122000000 + if not full and not ce_win and seen.get(va) == pa and n == 1: continue - todo.append((va, pa, n)) + todo[(va, pa, n)] = True + if full: + for va, pa, n in _coalesce_seen(seen): + todo[(va, pa, n)] = True + todo = list(todo) if not todo: return 0 @@ -202,6 +568,7 @@ def _flush(runs, seen): CHUNK = 200 for start in range(0, len(script), CHUNK): batch = script[start:start + CHUNK] + _heartbeat() try: r = subprocess.run(["sudo", "-n", "sh", "-c", "\n".join(batch)], capture_output=True, text=True, timeout=60) @@ -235,4 +602,12 @@ def _flush(runs, seen): if __name__ == "__main__": - main() + # main() loops forever; only an unexpected exception exits it. Restart + # instead of dying silently - a dead listener means no pins and every + # later CUDA client hangs at first cudaMalloc. + while True: + try: + main() + except Exception as e: + log("FATAL: main() exited: %r - restarting in 2s" % (e,)) + time.sleep(2) diff --git a/tools/guest-diag/pin_supervisor.sh b/tools/guest-diag/pin_supervisor.sh new file mode 100755 index 0000000..b87d5c1 --- /dev/null +++ b/tools/guest-diag/pin_supervisor.sh @@ -0,0 +1,169 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# Copyright (c) 2026 Open-IOV contributors +# pin_supervisor.sh - supervise the host-side pin tools (pin_gr_stream.py and +# semfix_uvmt.py) so a silent death cannot strand every later CUDA client at +# first cudaMalloc. +# +# Both tools write a heartbeat file; this script checks, every 5 s, that each +# process is alive (exact script path) and that its heartbeat mtime is < 75 s. +# A startup grace (120 s, skipped early once the tool log shows it has done +# first useful work) keeps a slow cold start from being restarted. On failure +# it kills the tool (exact path match) and restarts it, appending RESTART +# lines to the supervisor log. After 10 restarts in one hour for a tool it +# logs GIVEUP and exits 1. Only one semfix_uvmt instance is ever started +# (QEMU's QMP socket accepts a single client). +# +# Liveness policy (do not tighten without re-checking the soak evidence): +# * the tools refresh their heartbeats from dedicated loops/threads even +# while idle (pin_gr_stream: 5 s thread; semfix: per loop + around every +# blocking guest-agent/dmesg call), so 75 s never trips on a healthy +# client between phases; +# * a DEAD process is restarted on the next 5 s tick regardless of grace; +# * dead-pid/state recovery (the kill cascade) is handled INSIDE the tools +# (in-band #nopid + idle liveness probe + stall watchdog), NOT by this +# supervisor - a log-based restart here could fire mid-run of a healthy +# client (the m178 lesson). +# +# Usage: pin_supervisor.sh LOGDIR TAG +# Logs: LOGDIR/12-pin-supervisor-TAG.log (supervisor) +# LOGDIR/12-pin_gr_stream-TAG.log, LOGDIR/12-semfix-TAG.log (tools) +set -uo pipefail + +LOGDIR="${1:?usage: pin_supervisor.sh LOGDIR TAG}" +TAG="${2:?usage: pin_supervisor.sh LOGDIR TAG}" + +HERE=$(cd "$(dirname "$0")" && pwd) +PIN="$HERE/pin_gr_stream.py" +SEMFIX="$HERE/semfix_uvmt.py" +PIN_HB="${PIN_GR_HEARTBEAT:-/tmp/opencode/pin_gr_heartbeat}" +SEMFIX_HB="${SEMFIX_HEARTBEAT_FILE:-/tmp/opencode/semfix_heartbeat}" +SUPLOG="$LOGDIR/12-pin-supervisor-$TAG.log" +PIN_LOG="$LOGDIR/12-pin_gr_stream-$TAG.log" +SEMFIX_LOG="$LOGDIR/12-semfix-$TAG.log" +WINDOW=3600 +CAP=10 +STALE=75 +GRACE=120 +mkdir -p "$LOGDIR" /tmp/opencode + +log() { echo "[$(date +%H:%M:%S)] $*" | tee -a "$SUPLOG"; } + +tool_alive() { pgrep -f "python3 $1" >/dev/null 2>&1; } + +hb_age() { # age in seconds of the heartbeat file, 9999 if missing + local f="$1" m + m=$(stat -c %Y "$f" 2>/dev/null) || { echo 9999; return; } + echo $(( $(date +%s) - m )) +} + +kill_tool() { + pkill -f "python3 $1" 2>/dev/null + sleep 1 + if pgrep -f "python3 $1" >/dev/null 2>&1; then + pkill -9 -f "python3 $1" 2>/dev/null + sleep 1 + fi + return 0 +} + +start_pin() { + if tool_alive "$PIN"; then return 0; fi + touch "$PIN_HB" 2>/dev/null || true + ( setsid nohup env SEMFIX_COUNTER_RETRY=10 \ + python3 "$PIN" >> "$PIN_LOG" 2>&1 & ) + pin_started=$(date +%s) + log "START pin_gr_stream" +} + +start_semfix() { + # pgrep exact-count check before start: if an instance already exists + # (ours or a leftover), never start a second one. + if tool_alive "$SEMFIX"; then + log "semfix already running - not starting another" + return 0 + fi + touch "$SEMFIX_HB" 2>/dev/null || true + ( setsid nohup env SEMFIX_SKIP_PB=0 SEMFIX_COUNTER_RETRY=10 \ + python3 "$SEMFIX" 0.5 >> "$SEMFIX_LOG" 2>&1 & ) + semfix_started=$(date +%s) + log "START semfix_uvmt" +} + +prune() { # $1 now, $2 list -> list with entries older than WINDOW dropped + local now="$1" t out="" + for t in $2; do + [ $(( now - t )) -lt "$WINDOW" ] && out="$out $t" + done + printf '%s' "$out" +} + +count_list() { # $1 list -> element count + local n=0 t + for t in $1; do n=$(( n + 1 )); done + echo "$n" +} + +restarts_pin="" +restarts_semfix="" +pin_started=$(date +%s) +semfix_started=$(date +%s) + +restart() { # $1 label, $2 script path, $3 start function + local now list n + now=$(date +%s) + if [ "$1" = pin_gr_stream ]; then + restarts_pin=$(prune "$now" "$restarts_pin") + n=$(count_list "$restarts_pin") + else + restarts_semfix=$(prune "$now" "$restarts_semfix") + n=$(count_list "$restarts_semfix") + fi + if [ "$n" -ge "$CAP" ]; then + log "GIVEUP $1: $n restarts within ${WINDOW}s" + exit 1 + fi + kill_tool "$2" + if [ "$1" = pin_gr_stream ]; then + restarts_pin="$restarts_pin $now" + else + restarts_semfix="$restarts_semfix $now" + fi + log "RESTART $(( n + 1 )) $1 $(date +%Y-%m-%dT%H:%M:%S)" + "$3" +} + +check_tool() { # $1 label, $2 path, $3 heartbeat file, $4 start function, + # $5 tool start TS, $6 tool log file, $7 first-work pattern + local age runtime + age=$(hb_age "$3") + if tool_alive "$2" && [ "$age" -le "$STALE" ]; then + return 0 + fi + if tool_alive "$2"; then + # Startup grace: a cold tool can spend its first ~2 min installing the + # guest scanner or connecting QMP with a lagging heartbeat; only + # enforce staleness once the grace has passed or the log shows the + # tool has already done successful work. + runtime=$(( $(date +%s) - $5 )) + if [ "$runtime" -lt "$GRACE" ] && \ + ! grep -q -- "$7" "$6" 2>/dev/null; then + return 0 + fi + log "$1 heartbeat stale (${age}s) - restarting" + else + log "$1 not running (heartbeat age ${age}s) - restarting" + fi + restart "$1" "$2" "$4" +} + +log "pin_supervisor up (LOGDIR=$LOGDIR TAG=$TAG)" +start_pin +start_semfix +while true; do + sleep 5 + check_tool pin_gr_stream "$PIN" "$PIN_HB" start_pin "$pin_started" \ + "$PIN_LOG" "listening on vsock" + check_tool semfix_uvmt "$SEMFIX" "$SEMFIX_HB" start_semfix \ + "$semfix_started" "$SEMFIX_LOG" "semfix_uvmt: watching" +done diff --git a/tools/guest-diag/semfix_uvmt.py b/tools/guest-diag/semfix_uvmt.py index cf54445..11c191c 100755 --- a/tools/guest-diag/semfix_uvmt.py +++ b/tools/guest-diag/semfix_uvmt.py @@ -53,14 +53,47 @@ SKIP_PB = os.environ.get("SEMFIX_SKIP_PB", "0") == "1" # Seconds between counter-page pin attempts (the guest-agent exec blocks for # its full timeout when the client saturates the guest). COUNTER_RETRY_S = float(os.environ.get("SEMFIX_COUNTER_RETRY", "10")) +# Liveness file checked by pin_supervisor.sh (mtime must stay < 75 s; the +# supervisor grants a 120 s startup grace on top). +SEMFIX_HB_FILE = os.environ.get("SEMFIX_HEARTBEAT_FILE", + "/tmp/opencode/semfix_heartbeat") +# Dead-pid/stall watchdog tuning (seconds). The tracked client is probed in +# the guest only when it has been idle this long, on a cadence, so live +# clients are never probed and a killed client is detected within ~1 min. +LIVENESS_IDLE_S = float(os.environ.get("SEMFIX_LIVENESS_IDLE", "15")) +LIVENESS_EVERY_S = float(os.environ.get("SEMFIX_LIVENESS_EVERY", "25")) +# Pre-capture stall: client alive, counter page not pinned, no progress for +# this long -> recovery attempt; escalate loudly if it persists. +STALL_PRECAP_S = float(os.environ.get("SEMFIX_STALL_PRECAP", "60")) +STALL_RETRY_S = float(os.environ.get("SEMFIX_STALL_RETRY", "30")) +STALL_ESCALATE_S = float(os.environ.get("SEMFIX_STALL_ESCALATE", "180")) +# Inter-tool re-assert request: written here, consumed by pin_gr_stream.py +# (full ping re-assert of every pair it has ever pinned). +REASSERT_REQ = os.environ.get("PIN_REASSERT_FILE", + "/tmp/opencode/pin_reassert_request") + + +def _hb(): + try: + with open(SEMFIX_HB_FILE, "w") as f: + f.write("%.3f\n" % time.time()) + except OSError: + pass class Stream: """Tail the guest scanner's control lines mirrored by pin_gr_stream.py. - The guest emits "#pid <n>" and "#sem 0x<k>" over vsock; the listener - appends them here. This replaces per-loop guest-agent trace reads, which - were too slow to catch short-lived tracking semaphores. + The guest emits "#pid <n> [starttime]" and "#sem 0x<k>" over vsock; the + listener appends them here. This replaces per-loop guest-agent trace + reads, which were too slow to catch short-lived tracking semaphores. + + Replay safety (the m178 regression): the file persists across tool + restarts, so a fresh process re-reads the whole history. Lines written + before this process started (byte offset < init_size) are absorbed as + HISTORY - they seed the dedupe/identity state and the sem/pb lists but + can never trigger a per-client reset. Only lines appended afterwards + are live. This makes a supervisor restart mid-run harmless. """ def __init__(self, path): @@ -70,6 +103,15 @@ class Stream: self.pids = [] self.pbs = [] self.new_pbs = [] + self.new_nopids = [] + self.history_pids = [] # (pid, start) identities from history + self.pid_keys = set() # (pid, start) dedupe, per generation + try: + # Everything already in the file when this process starts is + # history (a restart replays it); never reset on it. + self._init_size = os.path.getsize(path) + except OSError: + self._init_size = 0 def reset(self): """Forget everything learned about the previous client. A new client @@ -80,6 +122,16 @@ class Stream: self.pbs = [] self.pids = [] self.new_pbs = [] + self.new_nopids = [] + self.pid_keys = set() + + @staticmethod + def _ident(line): + """Parse '#pid 1234' or '#pid 1234 5678900' -> (pid, start|None).""" + nums = re.findall(r"\d+", line) + if not nums: + return None, None + return int(nums[0]), (int(nums[1]) if len(nums) > 1 else None) def poll(self): try: @@ -90,12 +142,19 @@ class Stream: self.off = 0 f.seek(self.off) data = f.read() + base = self.off self.off = f.tell() except OSError: - return [], [] + return [], [], [] new_sems, new_pids = [], [] self.new_pbs = [] + self.new_nopids = [] + pos = base for line in data.splitlines(): + # A line is live when it ends past the startup boundary; history + # lines are absorbed (identity/dedupe) but never emitted as live. + live = pos + len(line) + 1 > self._init_size + pos += len(line) + 1 if line.startswith("#sem"): m = re.search(r"0x([0-9a-fA-F]+)", line) if m: @@ -111,13 +170,116 @@ class Stream: self.pbs.append(v) self.new_pbs.append(v) elif line.startswith("#pid"): - m = re.search(r"(\d+)", line) - if m: - v = int(m.group(1)) - if v not in self.pids: - self.pids.append(v) - new_pids.append(v) - return new_sems, new_pids + pid, start = self._ident(line) + if pid is not None and (pid, start) not in self.pid_keys: + self.pid_keys.add((pid, start)) + self.pids.append(pid) + if live: + new_pids.append((pid, start)) + else: + self.history_pids.append((pid, start)) + elif line.startswith("#nopid"): + # In-band client-death signal (new guest scanner; older ones + # simply never send it and the host probe covers the case). + pid, start = self._ident(line) + if live and pid is not None: + self.new_nopids.append((pid, start)) + return new_sems, new_pids, self.new_nopids + + +class ClientTracker: + """Monotonic, replay-proof client-identity tracking. + + A client's identity is (pid, starttime). Rules, in priority order: + * an identity adopted once is never reset a second time (duplicate + lines, trace/stream replays, supervisor restarts); + * a backwards pid, or the same pid with an older starttime, is a + replay - never reset; + * a pid recycled with a NEWER starttime is a real new client - reset + immediately (the previous holder is definitionally gone); + * a newer unseen identity only resets the per-client state once the + CURRENT client is CONFIRMED gone - a reset fired while the previous + client was still active is the m178 third-client hang regression. + + Pure logic: the caller feeds guest liveness results in and performs the + actual state purge when an action says so. + """ + + def __init__(self): + self.seen = set() + self.current = None + self.pending = None + self.dead = set() + + @staticmethod + def _start_match(a, b): + # Unknown starttimes match conservatively (legacy scanner lines). + return a[1] == b[1] or a[1] is None or b[1] is None + + def observe_history(self, pid, start): + """Absorb an identity found in pre-restart history: the process has + no per-client state yet, so history only seeds identity tracking and + the freshest identity becomes current - no reset is ever requested + for history lines.""" + key = (pid, start) + if key in self.seen: + return + self.seen.add(key) + self.current = key + + def note_dead(self, pid, start): + """A #nopid or a successful dead probe: identity is gone.""" + key = (pid, start) + self.dead.add(key) + if self.current and self.current[0] == pid \ + and self._start_match(self.current, key): + self.current = None + if self.pending and self.pending[0] == pid \ + and self._start_match(self.pending, key): + self.pending = None + + def is_current(self, pid, start): + """True when (pid, start) is the currently tracked identity (used to + decide whether a death signal may purge - a stale #nopid for a + recycled pid must never purge the new holder).""" + return self.current is not None and self.current[0] == pid \ + and self._start_match(self.current, (pid, start)) + + def decide(self, pid, start): + """Classify a live #pid line: adopt | skip | skip_backwards | + maybe_reset | reset.""" + key = (pid, start) + if key in self.seen: + return "skip" + self.seen.add(key) + cur = self.current + if cur is None: + self.current = key + return "adopt" + if pid == cur[0] and self._start_match(cur, key): + return "skip" + if pid == cur[0] and start is not None and cur[1] is not None \ + and start > cur[1]: + # Recycled pid, newer process: safe to reset now. + self.current = key + return "reset" + if pid < cur[0] or (pid == cur[0] and start is not None + and cur[1] is not None and start < cur[1]): + return "skip_backwards" + self.pending = key + return "maybe_reset" + + def resolve(self, old_dead): + """Caller learned whether the CURRENT client is dead. Returns + 'reset' (purge + adopt the pending client) or 'defer'.""" + if self.pending is None: + return None + if old_dead: + key = self.pending + self.pending = None + self.current = key + return "reset" + return "defer" def log(msg): @@ -196,7 +358,9 @@ def dmesg_releases(seen): Only lines after the last "sema-svc: thread started" marker count: dmesg keeps the previous module load's lines and the old client's releases. + Returns the number of newly printed lines (used as watchdog progress). """ + n = 0 try: r = subprocess.run( ["sudo", "-n", "sh", "-c", @@ -205,13 +369,15 @@ def dmesg_releases(seen): "| grep 'sema-push' | tail -300"], capture_output=True, text=True, timeout=10) except Exception: - return + return 0 for line in r.stdout.splitlines(): line = line.strip() if not line or line in seen: continue seen.add(line) + n += 1 log("push: %s" % line.replace("vgsp: ", "")) + return n def cdt_pid(): @@ -224,6 +390,34 @@ def cdt_pid(): return int(pids[-1]) if pids else None +def _pid_alive_cmd(pid): + """Guest shell that reports ALIVE / DEAD for a pid. Zombie-aware: a + killed (SIGKILLed) client lingers as a defunct /proc entry, and a plain + `test -d /proc/<pid>` would latch it as alive forever (the m178 + dead-pid latch).""" + return ("f=/proc/%d/stat; if [ -r \"$f\" ]; then " + "s=$(sed -e 's/^[^)]*) //' -e 's/ .*//' \"$f\"); " + "if [ \"$s\" = Z ] || [ -z \"$s\" ]; then echo DEAD; " + "else echo ALIVE; fi; else echo DEAD; fi" % pid) + + +def pid_alive(pid): + """Guest liveness of a client pid: True alive, False dead/zombie/gone, + None unknown (agent error). Callers must treat None as alive (never + purge a client on an inconclusive probe).""" + try: + r = subprocess.run([VMEXEC, "exec", _pid_alive_cmd(pid)], + capture_output=True, text=True, timeout=20) + except Exception: + return None + out = r.stdout or "" + if "ALIVE" in out: + return True + if "DEAD" in out: + return False + return None + + def pin_counter_page(pid, pinned_pids): """Pin the client's completion-counter page from its own pagemap. @@ -268,10 +462,87 @@ def main(): dseen = set() counter_pids = set() counter_next_try = {} + # Current client pid/starttime, mirrored from the tracker (starttime + # disambiguates recycled pids). last_pid = None + last_start = None + # Stale-sem tracking: consecutive unreadable reads, first-unreadable time, + # and the per-sem read skip deadline once a sem has failed 3 times. + unread_count = {} + first_unreadable_t = {} + skip_until = {} + # Failed pool-pin tracking: retry with a 15 s backoff after 2 failures. + sem_pin_fail = {} + sem_pin_skip_until = {} + qmp_next_try = 0.0 stream = Stream(SEMSTREAM) + tracker = ClientTracker() + # Dead-pid/backwards-transition proofing and the stall watchdog. + history_absorbed = False + adopted_from_history = False + probe_next = 0.0 # next allowed idle-liveness probe + recover_next = 0.0 # next watchdog recovery attempt + stall_since = None # first observation of the pre-capture stall + escalated = False + progress_t = time.time() # last host-side progress for the client loop = 0 + def _clear_client_state(): + """Purge every per-client latch. ONLY valid when the previous client + is CONFIRMED dead (or at fresh process start): the m178 third-client + hang came from wiping the state of an ACTIVE client on a replayed + (#pid) identity.""" + stream.reset() + known_sems.clear() + seen_sem_pin.clear() + last_state.clear() + last_seen_t.clear() + pinned_ranges.clear() + seen_ch.clear() + seen_pb.clear() + counter_pids.clear() + counter_next_try.clear() + sem_pin_fail.clear() + sem_pin_skip_until.clear() + unread_count.clear() + first_unreadable_t.clear() + skip_until.clear() + + def _adopt(pid, start, why): + nonlocal last_pid, last_start, progress_t, adopted_from_history + nonlocal stall_since, escalated + last_pid, last_start = pid, start + progress_t = time.time() + adopted_from_history = False + stall_since = None + escalated = False + log("semfix: tracking client pid %d%s (%s)" + % (pid, " start=%s" % start if start is not None else "", why)) + + def _watchdog_reset(): + nonlocal stall_since, escalated + stall_since = None + escalated = False + + def _watchdog_recover(): + """Idempotent pre-capture stall recovery: re-arm every sem latch so + the next loop re-resolves and re-pins, force an immediate counter + retry, and request a full pin re-assert from pin_gr_stream.""" + for sem in set(known_sems) | set(stream.sems): + skip_until.pop(sem, None) + sem_pin_skip_until.pop(sem, None) + sem_pin_fail.pop(sem, None) + unread_count.pop(sem, None) + first_unreadable_t.pop(sem, None) + seen_sem_pin.discard(sem) + if last_pid is not None: + counter_next_try[last_pid] = 0.0 + try: + with open(REASSERT_REQ, "w") as f: + f.write("%.3f\n" % time.time()) + except OSError: + pass + def pin_range_once(gpu_vas, cpu, gpa, size=0): key = tuple(gpu_vas), cpu, size st = pinned_ranges.get(key) @@ -293,40 +564,104 @@ def main(): while True: loop += 1 + _hb() sems = set() pbs = set() chans = set() - new_sems, new_pids = stream.poll() + new_sems, new_pids, new_nopids = stream.poll() for s in new_sems: log("stream: new uvmt sem 0x%x" % s) - for p in new_pids: - log("stream: client pid %d" % p) - if last_pid is not None and p != last_pid: - stream.reset() - known_sems.clear() - seen_sem_pin.clear() - last_state.clear() - last_seen_t.clear() - pinned_ranges.clear() - seen_ch.clear() - seen_pb.clear() - counter_pids.clear() - log("semfix: new client pid %d - reset per-client pin state" % p) - last_pid = p + if new_sems and last_pid is not None: + progress_t = time.time() + + # In-band client-death signals (new guest scanner): purge the dead + # client's per-client state immediately so the NEXT client works + # without a VM cycle (the kill cascade). + for p, st in new_nopids: + log("stream: client gone (%d%s)" + % (p, " start=%s" % st if st is not None else "")) + was_current = tracker.is_current(p, st) + tracker.note_dead(p, st) + if was_current and last_pid == p: + _clear_client_state() + last_pid = last_start = None + _watchdog_reset() + log("semfix: purged per-client state of exited pid %d" % p) + # The old client is gone: adopt any client whose reset was deferred + # while it was still (assumed) alive. + if last_pid is None and tracker.pending is not None \ + and tracker.current is None: + if tracker.resolve(old_dead=True) == "reset": + p, st = tracker.current + _clear_client_state() + _adopt(p, st, "deferred reset") + # Pre-restart history (a supervisor restart re-reads the file): absorb + # identities WITHOUT any reset - the fresh process already has empty + # state; the freshest identity is the client that was running. + if stream.history_pids and not history_absorbed: + for p, st in stream.history_pids: + tracker.observe_history(p, st) + history_absorbed = True + if tracker.current is not None and last_pid is None: + p, st = tracker.current + _adopt(p, st, "resumed from pre-restart history") + adopted_from_history = True + + for p, st in new_pids: + cur_before = tracker.current + action = tracker.decide(p, st) + tag = " start=%s" % st if st is not None else "" + if action == "skip": + log("stream: client pid %d%s seen before - no reset" % (p, tag)) + elif action == "skip_backwards": + log("semfix: client pid %d%s backwards/replay - no reset" + % (p, tag)) + elif action == "adopt": + log("stream: client pid %d%s" % (p, tag)) + _adopt(p, st, "new client") + elif action == "reset": + # Recycled pid, newer starttime: the previous holder of this + # pid is definitionally gone, safe to reset now. + log("semfix: pid %d recycled (start=%s) - reset per-client " + "pin state" % (p, st)) + _clear_client_state() + _adopt(p, st, "pid recycled") + elif action == "maybe_reset": + old = cur_before + alive = pid_alive(old[0]) if old else False + if alive is False: + log("semfix: new client pid %d%s (pid %d gone) - reset " + "per-client pin state" + % (p, tag, old[0] if old else -1)) + _clear_client_state() + tracker.resolve(old_dead=True) + _adopt(p, st, "client switch") + else: + # NEVER clear the active client's state on the account of + # a newer pid: the m178 hang. Deferred until the old pid + # is confirmed gone (periodic idle probe below). + log("semfix: client pid %d%s appeared while pid %s is " + "alive%s - reset deferred" + % (p, tag, old[0] if old else "?", + "" if alive else " (probe unknown)")) for pb in getattr(stream, "new_pbs", []): log("stream: new pushbuffer 0x%x" % pb) pbs.add(pb) + if getattr(stream, "new_pbs", []) and last_pid is not None: + progress_t = time.time() sems.update(stream.sems) # The guest-agent trace read is slow under client load; use it only # as a fallback for channel/pushbuffer discovery and to recover sems # the scanner may have missed. if loop % TRACE_EVERY == 0: + _hb() try: t_sems, t_pbs, t_chans = guest_trace() except Exception: t_sems, t_pbs, t_chans = set(), set(), set() + _hb() sems.update(t_sems) pbs.update(t_pbs) chans.update(t_chans) @@ -336,16 +671,54 @@ def main(): for sem in sorted(sems - known_sems): known_sems.add(sem) - for sem in sorted(known_sems): + for hb_i, sem in enumerate(sorted(known_sems)): + if hb_i % 16 == 0: + # A full pass can run 60-90 s (thousands of QMP reads); keep + # the heartbeat fresh inside it (r203: 79 s gap -> supervisor + # restarted the tool mid-pass). + _hb() + now = time.time() + if now < skip_until.get(sem, 0): + continue try: st = uvmt_state(q, sem) except Exception: st = None + if st is None: + # Stale-sem hygiene: the trace ring replays dead clients' + # kernel pointers forever; without eviction every loop burns + # QMP reads on them and the log floods with "unreadable". + cnt = unread_count.get(sem, 0) + 1 + unread_count[sem] = cnt + first_unreadable_t.setdefault(sem, now) + if cnt >= 3: + skip_until[sem] = now + 30.0 + if cnt == 3: + log("semfix: sem=0x%x unreadable x3 - retrying every " + "30s" % sem) + age = now - first_unreadable_t[sem] + if age >= 60.0 and sem not in stream.sems: + known_sems.discard(sem) + unread_count.pop(sem, None) + first_unreadable_t.pop(sem, None) + skip_until.pop(sem, None) + last_state.pop(sem, None) + last_seen_t.pop(sem, None) + seen_sem_pin.discard(sem) + sem_pin_fail.pop(sem, None) + sem_pin_skip_until.pop(sem, None) + log("semfix: evicting stale sem=0x%x (unreadable %.0fs, " + "absent from stream)" % (sem, age)) + continue + else: + unread_count.pop(sem, None) + first_unreadable_t.pop(sem, None) + skip_until.pop(sem, None) key = state_key(st) - now = time.time() if key != last_state.get(sem): last_state[sem] = key last_seen_t[sem] = now + progress_t = time.time() log("uvmt %s" % state_str(sem, st)) elif now - last_seen_t.get(sem, 0) >= HEARTBEAT: last_seen_t[sem] = now @@ -354,18 +727,34 @@ def main(): # Pin the pool page once per (sem, word) resolution. The payload # word is what the releasing push writes; pinning the page makes # the host able to apply that release. - if st and sem not in seen_sem_pin: + if (st and sem not in seen_sem_pin + and now >= sem_pin_skip_until.get(sem, 0)): try: pairs = resolve_sem_pool(q, sem) except Exception as e: pairs = [] log("semfix: sem=0x%x resolve failed: %s" % (sem, e)) if pairs: - seen_sem_pin.add(sem) + ok_any = False for gv, gpa in pairs: ok = pin(gv, gpa) + ok_any = ok_any or ok log("semfix: sem=0x%x pool va=0x%x -> gpa=0x%x pin=%s" % (sem, gv, gpa, "ok" if ok else "FAILED")) + if not ok_any: + # Mark pinned only after a successful pin(): a failed + # pin used to be recorded as done and never retried. + fails = sem_pin_fail.get(sem, 0) + 1 + sem_pin_fail[sem] = fails + if fails >= 2: + sem_pin_skip_until[sem] = now + 15.0 + log("semfix: sem=0x%x pool pin failed x%d - retry " + "in 15s" % (sem, fails)) + continue + seen_sem_pin.add(sem) + sem_pin_fail.pop(sem, None) + sem_pin_skip_until.pop(sem, None) + progress_t = time.time() if SKIP_PB: continue d = {} @@ -385,6 +774,7 @@ def main(): d["exception"] = str(e)[:40] if res: pin_range_once(*res) + progress_t = time.time() log("semfix: sem=0x%x pushbuf -> %s" % (sem, diag_str(d))) else: @@ -393,7 +783,22 @@ def main(): else: log("semfix: sem=0x%x pool unresolved (retry)" % sem) + if getattr(q, "_fails", 0) >= 5: + # QMP went away (QEMU restart, monitor died): all reads have been + # empty for 5 consecutive replies. Recreate the client loudly + # instead of reading empty forever ("uvmt(hb) ... unreadable"). + now = time.time() + if now >= qmp_next_try: + log("semfix: QMP unreadable (%d consecutive empty replies) - " + "recreating connection" % q._fails) + try: + q = Qmp(QMP_SOCK) + except Exception as e: + log("semfix: QMP recreate failed: %s" % e) + qmp_next_try = now + 5.0 + for ch in sorted(chans - seen_ch): + _hb() # a resolve can block for tens of seconds d = {} try: res = resolve_from_channel(q, ch, d) @@ -403,10 +808,12 @@ def main(): if res: seen_ch.add(ch) pin_range_once(*res) + progress_t = time.time() log("semfix: ch=0x%x -> %s" % (ch, diag_str(d))) else: log("semfix: ch=0x%x pushbuf miss [%s]" % (ch, diag_str(d))) for pb in sorted(pbs - seen_pb): + _hb() # a resolve can block for tens of seconds d = {} try: res = resolve_from_pushbuffer(q, pb, d) @@ -416,12 +823,15 @@ def main(): if res: seen_pb.add(pb) pin_range_once(*res) + progress_t = time.time() log("semfix: pb=0x%x -> %s" % (pb, diag_str(d))) else: log("semfix: pb=0x%x pushbuf miss [%s]" % (pb, diag_str(d))) if DLOG: - dmesg_releases(dseen) + if dmesg_releases(dseen): + progress_t = time.time() + _hb() # Retry incomplete push-buffer pins: gva2gpa of the kernel vmalloc # alias fails while the sampled vCPU runs in user mode (PTI), so a @@ -430,24 +840,89 @@ def main(): if loop % 8 == 0: for st in list(pinned_ranges.values()): if not st["done"] and st["tries"] > 0: + _hb() # pin_range_once can walk thousands of QMP reads pin_range_once(*st["args"]) + # Dead-pid detection and the pre-capture stall watchdog. The guest is + # probed ONLY when the tracked client has been quiet for a while (a + # live, pushing client is never probed), so the agent is not burdened + # during healthy runs. A confirmed-dead client is purged immediately: + # the NEXT client then starts with clean state - no VM cycle needed + # (the kill cascade). + now = time.time() + if adopted_from_history and last_pid is not None: + # The client resumed from pre-restart history may have died while + # this process was down; do not latch a dead pid into the counter + # retries. + alive = pid_alive(last_pid) + adopted_from_history = False + if alive is False: + log("semfix: resumed pid %d already gone - dropping" % last_pid) + tracker.note_dead(last_pid, last_start) + last_pid = last_start = None + _watchdog_reset() + if last_pid is not None and now >= probe_next \ + and now - progress_t >= LIVENESS_IDLE_S: + probe_next = now + LIVENESS_EVERY_S + alive = pid_alive(last_pid) + if alive is False: + log("semfix: client pid %d gone (probe) - purging per-client " + "state" % last_pid) + tracker.note_dead(last_pid, last_start) + _clear_client_state() + last_pid = last_start = None + _watchdog_reset() + if tracker.pending is not None and tracker.current is None \ + and tracker.resolve(old_dead=True) == "reset": + p, st = tracker.current + log("semfix: adopting deferred client pid %d" % p) + _adopt(p, st, "deferred reset") + elif alive is True and last_pid not in counter_pids: + # Pre-capture stall: the client is alive but has no pinned + # counter page and nothing progressed. Recover (idempotent), + # then escalate loudly if the stall persists - the goal is + # that a first cudaMalloc / first GR capture never stalls + # silently. + idle = now - progress_t + if idle >= STALL_PRECAP_S: + if stall_since is None: + stall_since = now + if now >= recover_next: + recover_next = now + STALL_RETRY_S + log("STALL WATCHDOG: client pid %d alive, no " + "pre-capture progress for %.0fs - recovery " + "attempt (re-arm sems, re-pin counter, request " + "pin re-assert)" % (last_pid, idle)) + _watchdog_recover() + if not escalated and now - stall_since >= STALL_ESCALATE_S: + escalated = True + log("STALL WATCHDOG: recovery FAILED for pid %d " + "after %.0fs (alive, counter page not pinned) - " + "pre-capture hang persists, needs manual " + "attention" % (last_pid, now - stall_since)) + # Pin the completion-counter page of a newly started client, before its # first GR pushes decode their releases. The pid comes from the guest # scanner; the guest-agent fallback is only a safety net. The exec can # block for its whole timeout under client load, so retry on a slow # cadence: hammering the agent here starves it and wedges the run. if last_pid is None and loop % TRACE_EVERY == 0: - last_pid = cdt_pid() - if last_pid: - log("counter: client pid %d (agent)" % last_pid) + p = cdt_pid() + _hb() + if p: + if tracker.current is None and (p, None) not in tracker.seen: + tracker.observe_history(p, None) + last_pid, last_start = p, None + log("counter: client pid %d (agent)" % p) if last_pid and last_pid not in counter_pids: now = time.time() if now >= counter_next_try.get(last_pid, 0): try: - pin_counter_page(last_pid, counter_pids) + if pin_counter_page(last_pid, counter_pids): + progress_t = time.time() except Exception as e: log("counter: pin failed: %s" % e) + _hb() counter_next_try[last_pid] = now + COUNTER_RETRY_S time.sleep(interval) diff --git a/tools/guest-diag/semfix_watch.py b/tools/guest-diag/semfix_watch.py index 5124805..4b60fd6 100644 --- a/tools/guest-diag/semfix_watch.py +++ b/tools/guest-diag/semfix_watch.py @@ -58,12 +58,47 @@ CHAN_MAX_CANDS = int(os.environ.get("CHAN_MAX_CANDS", "256")) class Qmp: def __init__(self, path): + self.path = path + self._cache = {} + self._fails = 0 + self._last_reconnect = 0.0 + self._reconnecting = False self.s = socket.socket(socket.AF_UNIX, socket.SOCK_STREAM) self.s.connect(path) self.s.settimeout(15) self._rd() self.cmd("qmp_capabilities") - self._cache = {} + + def reconnect(self): + """Re-establish the QMP connection after the socket died. + + QEMU's QMP socket accepts a single client and closes when the guest + reboots; cmd() used to swallow every exception and return {}, so all + reads stayed empty forever ("uvmt(hb) sem=... unreadable") and + pinning stopped. Rate-limited to one attempt per second so a failing + monitor does not turn every read into a connect storm.""" + if self._reconnecting or time.time() - self._last_reconnect < 1.0: + return + self._reconnecting = True + self._last_reconnect = time.time() + try: + try: + self.s.close() + except Exception: + pass + self.s = socket.socket(socket.AF_UNIX, socket.SOCK_STREAM) + self.s.connect(self.path) + self.s.settimeout(15) + self._rd() + self.cmd("qmp_capabilities") + self._cache = {} + self._ncpu = 0 + self._cpu = 0 + self._fails = 0 + except Exception: + pass + finally: + self._reconnecting = False def _rd(self): """Read one QMP reply (JSON object ends with a newline). Poll with a @@ -91,8 +126,17 @@ class Qmp: def cmd(self, name, **args): try: self.s.sendall(json.dumps({"execute": name, "arguments": args}).encode() + b"\n") - return json.loads(self._rd() or "{}") + reply = self._rd() + if not reply: + raise OSError("empty QMP reply") + self._fails = 0 + return json.loads(reply) except Exception: + # Broken socket (QEMU restart, monitor hiccup): try exactly one + # reconnect, rate-limited inside reconnect(). + self._fails = getattr(self, "_fails", 0) + 1 + if not getattr(self, "_reconnecting", False): + self.reconnect() return {} def hmp(self, line): @@ -143,8 +187,14 @@ class Qmp: break m = re.search(r":(.*)$", out, re.S) if not m: + if not out.strip(): + # Empty monitor output: the socket is broken, not the address. + self.reconnect() return [] vals = [int(x, 16) for x in re.findall(r"0x([0-9a-fA-F]+)", m.group(1))][:n] + if not vals: + self.reconnect() + return [] if cache: self._cache[key] = vals return vals @@ -189,18 +239,46 @@ def guest_trace(): Read *before* clearing: the create-time probes fire once per client and would be wiped by a clear that races them. + + 2026-10-16 robustness fix (the cascade root): this used to run a bare + `tail` over the guest agent. A tracefs read that blocks under a busy + guest left *uninterruptible* tail processes behind (observed dozens in + D state), the guest-agent exec then blocked for its full 90 s timeout + every iteration, and the pool/tracker pins this tool supplies stopped + flowing - after ~5 clients every new client froze at DEVICE waiting on + its UVM tracker with `pending release starved ... pool/waiter page`. + Now the read is bounded remotely (`timeout -s KILL 12`) so no tail can + outlive its call, and SSH (which stays healthy when the guest agent + wedges) is preferred with the agent as fallback. """ - inner = ("tail -c 131072 /sys/kernel/tracing/trace " + inner = ("timeout -s KILL 12 tail -c 131072 /sys/kernel/tracing/trace " "| grep -E \"uvmt|uvch|uvbp|uvpbf\" " "| tail -400") if not os.environ.get("KEEP_TRACE"): inner += "; echo > /sys/kernel/tracing/trace" - cmd = "sudo sh -c '%s'" % inner + outer = "sudo sh -c '%s'" % inner + out = "" + ssh = os.environ.get("SEMFIX_SSH_PORT", "2222") + ssh_cmd = [ + "sshpass", "-p", os.environ.get("VGPU_GUEST_PASS", "ubuntu"), + "ssh", "-o", "StrictHostKeyChecking=no", + "-o", "UserKnownHostsFile=/dev/null", + "-o", "PubkeyAuthentication=no", "-o", "ConnectTimeout=6", + "-p", ssh, "ubuntu@127.0.0.1", outer, + ] try: - out = subprocess.run([VMEXEC, "exec", cmd], capture_output=True, - text=True, timeout=90).stdout + r = subprocess.run(ssh_cmd, capture_output=True, text=True, + timeout=25) + out = r.stdout or "" except Exception: - return set(), set(), set() + out = "" + if not out: + try: + out = subprocess.run([VMEXEC, "exec", outer], + capture_output=True, text=True, + timeout=25).stdout or "" + except Exception: + return set(), set(), set() sems = {int(m, 16) for m in re.findall(r"sem=0x([0-9a-fA-F]+)", out)} pbs = {int(m, 16) for m in re.findall(r"pb=0x([0-9a-fA-F]+)", out)} chans = {int(m, 16) for m in re.findall(r"ch=0x([0-9a-fA-F]+)", out)} @@ -369,8 +447,11 @@ def resolve_pushbuffer_wide(q, sem, d): def pin(va, pa): cmd = "printf '0x%x,0x%x' > %s/sema_map_add" % (va, pa, PARAM) for _ in range(40): - r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], - capture_output=True, text=True) + try: + r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], + capture_output=True, text=True, timeout=5) + except subprocess.TimeoutExpired: + continue if r.returncode == 0 and "timed out" not in (r.stderr or ""): return True time.sleep(0.05) # async queue full: drain and retry @@ -380,8 +461,11 @@ def pin(va, pa): def pin_bulk(va, pa, nbytes): cmd = "printf '0x%x,0x%x,0x%x' > %s/sema_map_add" % (va, pa, nbytes, PARAM) for _ in range(40): - r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], - capture_output=True, text=True) + try: + r = subprocess.run(["sudo", "-n", "sh", "-c", cmd], + capture_output=True, text=True, timeout=5) + except subprocess.TimeoutExpired: + continue if r.returncode == 0 and "timed out" not in (r.stderr or ""): return True time.sleep(0.05) diff --git a/tools/lg/frame_reader.py b/tools/lg/frame_reader.py new file mode 100644 index 0000000..f8bab40 --- /dev/null +++ b/tools/lg/frame_reader.py @@ -0,0 +1,77 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: BSD-3-Clause +# frame_reader.py - read the GPU-rendered frames the guest publishes into the +# kvmfr ivshmem BAR (see tools/cuda-ab/frame_producer.cu) and save them on the +# host as PPM/PNG. This is the host half of the guest frame-producer path: +# the QEMU ivshmem-plain device is backed by /dev/shm/kvmfr-<uuid>, which the +# guest maps through /sys/bus/pci/devices/0000:00:03.0/resource2. +# +# Usage: frame_reader.py [--shm /dev/shm/kvmfr-<uuid>] [--frames N] [--out DIR] +import argparse +import mmap +import os +import struct +import sys +import time + +MAGIC = 0x4B564D46 +FRAME_OFFSET = 0x100000 +# magic, version, width, height, stride, format, seq, reserved, timestamp_ns +HDR = struct.Struct("<IIIIIIIIQ") + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--shm", default=None) + ap.add_argument("--frames", type=int, default=5) + ap.add_argument("--out", default="/tmp/opencode/frames") + ap.add_argument("--timeout", type=float, default=120.0) + args = ap.parse_args() + + if not args.shm: + import glob + cands = [p for p in glob.glob("/dev/shm/kvmfr-*") + if os.path.getsize(p) > 0] + if not cands: + print("NO_SHM_FILE", file=sys.stderr) + return 1 + args.shm = max(cands, key=os.path.getsize) + os.makedirs(args.out, exist_ok=True) + print("SHM %s" % args.shm) + + fd = os.open(args.shm, os.O_RDWR) + size = os.fstat(fd).st_size + mm = mmap.mmap(fd, size, mmap.MAP_SHARED, + mmap.PROT_READ | mmap.PROT_WRITE) + + deadline = time.time() + args.timeout + saved = 0 + last_seq = 0 + while saved < args.frames and time.time() < deadline: + mm.seek(FRAME_OFFSET) + magic, ver, w, h, stride, fmt, seq, _res, ts = HDR.unpack( + mm.read(HDR.size)) + if magic != MAGIC or seq == 0 or seq == last_seq or w == 0: + time.sleep(0.05) + continue + fb = mm.read(w * h * 4) + last_seq = seq + path = os.path.join(args.out, "frame-%04d.ppm" % seq) + with open(path, "wb") as f: + f.write(b"P6\n%d %d\n255\n" % (w, h)) + # RGBA -> RGB + rgb = bytearray(w * h * 3) + rgb[0::3] = fb[0::4] + rgb[1::3] = fb[1::4] + rgb[2::3] = fb[2::4] + f.write(bytes(rgb)) + saved += 1 + print("SAVED seq=%d %dx%d fmt=%d -> %s" % (seq, w, h, fmt, path)) + mm.close() + os.close(fd) + print("READER_DONE saved=%d last_seq=%d" % (saved, last_seq)) + return 0 if saved else 2 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tools/lg/frame_sequence.sh b/tools/lg/frame_sequence.sh new file mode 100755 index 0000000..958a34a --- /dev/null +++ b/tools/lg/frame_sequence.sh @@ -0,0 +1,33 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# frame_sequence.sh - capture an animated GPU-rendered frame sequence end to +# end: run the guest's CUDA frame producer into the kvmfr ivshmem BAR and +# save every published frame on the host as PPM/PNG (the proven +# CUDA -> ivshmem -> host path; see docs/05-operations/06). +# +# Usage: tools/lg/frame_sequence.sh [frames] [width] [height] [fps] [outdir] +set -u +HERE="$(cd "$(dirname "$0")" && pwd)" +REPO="$(cd "$HERE/../.." && pwd)" +FRAMES="${1:-40}" +W="${2:-800}" +H="${3:-600}" +FPS="${4:-10}" +OUT="${5:-/tmp/opencode/frames-seq}" + +mkdir -p "$OUT" +rm -f "$OUT"/frame-*.ppm "$OUT"/frame-*.png + +# Start the host reader first so no published frame is missed. +python3 "$REPO/tools/lg/frame_reader.py" --frames "$FRAMES" --out "$OUT" \ + --timeout 240 > "$OUT/reader.log" 2>&1 & +READER=$! + +sleep 1 +"$REPO/scripts/ssh_guest.sh" \ + "sudo /tmp/frame_producer $FRAMES $W $H $FPS" 2>&1 | grep -v "Permanently added" \ + | tail -5 + +wait "$READER" +echo "reader rc=$?" +ls -1 "$OUT"/frame-*.ppm 2>/dev/null | wc -l diff --git a/tools/site/pages/Main-Page.wiki b/tools/site/pages/Main-Page.wiki index 9a80009..f91fc17 100644 --- a/tools/site/pages/Main-Page.wiki +++ b/tools/site/pages/Main-Page.wiki @@ -18,7 +18,7 @@ The only new component is one host kernel module (<code>vgsp.ko</code>), which i * '''[[Documentation]]''' — the full index of this wiki. * '''[[FAQ]]''' — the questions we get asked most. -== Project status (2026-10-12) == +== Project status (2026-10-20) == {| class="wikitable" ! Area !! State |- @@ -30,9 +30,9 @@ The only new component is one host kernel module (<code>vgsp.ko</code>), which i |- | CUDA userspace init and the driver-API smoke (ALL OK, 64-byte payload round-trips) || Working |- -| Guest CUDA kernel execution on real hardware || In progress — a forwarded launch reaches the engine queue; engine execution is the frontier +| Guest CUDA kernel execution on real hardware || Working — vector-add, matrix-multiply and transformer kernels run green on fresh boots; hardening and breadth in progress |- -| OpenGL / Vulkan, Windows guests, multi-tenancy, migration, performance || Not yet +| OpenGL / Vulkan, Windows guests, multi-tenancy, migration, performance || Not yet — multiple mediated devices (different emulated VRAM sizes) and multiple guests already coexist; per-guest execution isolation is in progress |} The full story, with evidence and the current wall: [[Project Status]] and [[The Execution Frontier]]. |