diff options
| -rw-r--r-- | kernel/gsp_sema_service.c | 274 | ||||
| -rw-r--r-- | tools/cuda-ab/memset2d_probe.cu | 157 | ||||
| -rwxr-xr-x | tools/cuda-ab/run/memset2d_probe | bin | 0 -> 820904 bytes | |||
| -rwxr-xr-x | tools/cuda-ab/run/streams_events_probe | bin | 0 -> 829304 bytes | |||
| -rwxr-xr-x | tools/cuda-ab/run/tgemm_probe | bin | 0 -> 829192 bytes | |||
| -rw-r--r-- | tools/cuda-ab/streams_events_probe.cu | 170 | ||||
| -rw-r--r-- | tools/cuda-ab/tgemm_probe.cu | 216 | ||||
| -rw-r--r-- | tools/guest-diag/pin_gr_stream.py | 2 |
8 files changed, 789 insertions, 30 deletions
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index bcd8ac3..ea96595 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -420,6 +420,18 @@ struct svc_channel { u64 gr_nocap_va; /* push VA the window was opened for */ unsigned long gr_nocap_since; /* absolute first-seen stamp */ /* + * r471: the completion-counter publish must never advance the + * guest's word past a push whose content was still unwritten when + * the entry was consumed (the streams_events staging-reuse + * corruption: the publish satisfied the client's D2H wait before + * its copy ran, and the readback drained a staging buffer that + * still held the PREVIOUS copy's bytes). Lowest ring slot consumed + * with all-zero content; 0 = none. Cleared when a later decode of + * a GR push reads non-zero content (the content landed - the rewinds + * re-read entries constantly, so this heals within a pass or two). + */ + u32 gr_bad_put; + /* * STALLFREE r225: stale-epoch / stale-replay wedge watchdog. A * channel that cannot make forward progress (last_get pinned behind * put, or a residual session put that never drops) is detected here @@ -486,6 +498,12 @@ struct svc_vaspace { u64 pdb; u32 entries; /* root entry count */ int ap; /* SVC_APERTURE_* of the directory itself */ + unsigned long stamp; /* r459: last registration/use; LRU evict + * when the table is full (a stale entry from + * a dead client otherwise blocks every new + * client after the 8th registration: the + * mid-boot walk failure that made the 10th+ + * client's copies fail rc=-2) */ }; struct svc_lat { @@ -819,6 +837,7 @@ struct sema_svc { u64 pushtrace_all; /* fail4: de-windowed push decode * trace (every push, bounded) */ u64 opadd_log; /* fail4: every op appended */ + u64 op2d_raw_log; /* r458: raw 2D-shaped copy pushes */ u64 dmatrace_log; /* fail4: denied-transfer trace */ u64 pend_dbg_log; /* fail4 r305: pending/hold witness */ u64 snap_gate_log; /* fail4 r311: capture-gate witness */ @@ -1599,12 +1618,40 @@ void gsp_sema_service_vaspace_add(const struct sema_vaspace_desc *desc) if (!vs && !svc.vas[i].valid) vs = &svc.vas[i]; } + /* + * r459: the table is 8 slots and the guest driver does not always + * send a teardown for a client's VA space (observed: 23 distinct + * registrations, 0 drops in one boot). Once all slots hold a dead + * client's directory, a NEW client's registration found no slot and + * every walk of its VAs failed (copies rc=-2, launches reading + * stale pages - the 10th-client mid-boot degradation). Evict the + * least-recently-registered entry instead of refusing the new one; + * the old client is gone and its directory pointer is dead guest + * memory, so the eviction can only remove stale authority. + */ + if (!vs) { + unsigned long oldest = 0; + int oi = 0; + + for (i = 0; i < SVC_MAX_VASPACES; i++) { + if (!svc.vas[i].valid) + continue; + if (!oldest || time_before(svc.vas[i].stamp, oldest)) { + oldest = svc.vas[i].stamp; + oi = i; + } + } + vs = &svc.vas[oi]; + pr_info("sema-svc: VA space table full; evicting 0x%x (pdb=0x%llx) for 0x%x\n", + vs->h_vaspace, vs->pdb, desc->h_vaspace); + } if (vs) { vs->valid = true; vs->h_vaspace = desc->h_vaspace; vs->pdb = desc->pdb; vs->entries = desc->num_entries; vs->ap = desc->b_vidmem ? SVC_APERTURE_FB : SVC_APERTURE_SYS; + vs->stamp = jiffies; /* A new directory invalidates what we learned from the old one. */ memset(svc.lat, 0, sizeof(svc.lat)); } @@ -1932,6 +1979,26 @@ static int svc_walk_vaspaces(struct nv_vgpu_device *vgpu, u64 va, } spin_unlock_irqrestore(&svc.vas_lock, flags); + /* + * r459: try the NEWEST registration first. A dead client's stale + * directory (the table outlives clients; only 8 slots) can still + * resolve a VA - through guest memory the driver has since reused - + * to a wrong page, and the first success wins. Ordering by the + * registration stamp makes the live client's directory the + * authority and only falls through to older ones when it cannot + * answer. + */ + for (i = 1; i < n; i++) { + struct svc_vaspace k = local[i]; + int j = i - 1; + + while (j >= 0 && time_before(local[j].stamp, k.stamp)) { + local[j + 1] = local[j]; + j--; + } + local[j + 1] = k; + } + for (i = 0; i < n; i++) { int alt; @@ -5997,14 +6064,18 @@ static void svc_copy_snap_free(u64 dst, u64 src, size_t total, } /* - * r444: drop every snapshot still owned by one entry. A snapshot that was - * not consumed by the entry's own copy execution (the copy ran via the - * forwarded hardware path, or a post-execution re-decode re-captured the - * unchanged staging) has no future consumer once the entry's release is - * delivered - leaving it behind let a LATER entry whose copy shares the - * (dst,src,total) triple execute from the stale bytes. + * r456/r457: the snapshot lifecycle is capture-once -> free on the copy's + * own successful consumption (svc_copy) -> free on ring-slot reuse with a + * new push hash (svc_copy_snap_capture) -> free when a re-decode skips an + * already-done op (svc_run_ops) -> TTL. The old per-entry owner free was + * removed from every opdef_clear caller: an entry whose release was + * delivered can still be REPLAYED (the pin-batch rewind) with its copy + * still owed, and freeing at that point let the replay re-capture the + * driver's refilled ping-pong staging (the r457 mode-C b1 H2D wrote b2's + * bytes into db1). A stale snapshot can only ever serve its own + * (ch, idx, hash) owner, so the r444 cross-entry poisoning stays closed. */ -static void svc_copy_snap_free_owner(u32 owner_ch, u32 owner_idx) +static void __maybe_unused svc_copy_snap_free_owner(u32 owner_ch, u32 owner_idx) { u32 i; @@ -9903,6 +9974,23 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, break; case CE_M_LINE_LENGTH_IN: st.line_length = v0; + /* + * r458 THE memset2d root cause: the driver encodes + * LINE_LENGTH_IN (0x418) and the adjacent LINE_COUNT + * (0x41c) as ONE count=2 method group - observed raw + * push bytes `20028106 00001000 00000100` = header + * (subch 4, count 2, addr 0x418), line length 4096, + * line count 0x100 = 256. Reading only v0 left + * line_count = 0, so the 2D-copy builder below took + * the single-row path: one 4 KB op per push, and a + * 4 MB cudaMemcpy2D delivered exactly one page per + * 1 MB block (memset2d probe: per-MB ok=4096, + * bad=12.5M/12.58M). The GRCE variant below has + * always decoded the count=2 form; the CE path did + * not. + */ + if (count >= 2) + st.line_count = v1; break; /* r402: 2D geometry (cudaMemset2D / cudaMemcpy2D). */ case CE_M_LINE_COUNT: @@ -10210,6 +10298,39 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, m->spitch = spitch; } } + /* + * r458 memset2d forensics: the + * driver's cudaMemcpy2D pushes decode + * LINE_LENGTH=4096 with PITCH_IN/OUT + * present but LINE_COUNT absent + * (lc=0), so only one 4 KB row per + * push executes and a 4 MB copy + * delivers 4 pages. Dump the raw push + * (bounded) for every pitch-shaped + * copy so the method stream that + * carries the row count is on record. + */ + if (gsp_sema_dbg && + (st.have_pitch_out || + st.have_pitch_in) && + svc.op2d_raw_log < 12) { + u32 dl = len < 220 ? len : 220; + + svc.op2d_raw_log++; + pr_info("sema-op2d-raw: h=0x%x push=0x%llx len=%u off=0x%x line=%u lc=%u po=%u/%u pi=%u/%u out=0x%llx in=0x%llx bytes=%u dma=0x%08x\n", + ch->h_object, push_va, + len, off, + st.line_length, + st.line_count, + st.have_pitch_out, + st.pitch_out, + st.have_pitch_in, + st.pitch_in, + m->dst, m->src, + (u32)m->bytes, dma); + pr_info("sema-op2d-raw: raw %*phN\n", + dl, buf); + } n_ops++; ops_seq++; } else if (gsp_sema_dbg && svc.dma_skip_log < 32) { @@ -14988,6 +15109,25 @@ static void svc_counter_publish(struct nv_vgpu_device *vgpu, return; val = put; + /* + * r471: never advance the guest's completion word past a push whose + * content was still unwritten when its entry was consumed. The + * clamp position is the lowest all-zero GR slot (gr_bad_put = slot + * + 1; 0 = none). Without it the publish can satisfy a D2H wait + * before its copy ran and the client drains a reused staging buffer + * holding the previous copy's bytes (streams_events h_c[2] = c0). + * The clamp heals when the rewinds re-read the slot with real + * content; the real completions still arrive through the decoded + * releases. + */ + if (!ch->engine_type && ch->gr_bad_put && val >= ch->gr_bad_put) { + val = ch->gr_bad_put - 1; + if (gsp_sema_dbg && svc.counter_pub_dbg < 96) { + svc.counter_pub_dbg++; + pr_info("sema-diag: pub-clamp h=0x%x pa=0x%llx val=%u (bad=%u) cur=%u\n", + ch->h_object, pa, val, ch->gr_bad_put, cur); + } + } if (val == ch->counter_last) return; /* @@ -15907,28 +16047,28 @@ static void svc_opdef_clear(struct svc_opdef_ent *e, u32 caller) e->n ? e->rel[0].payload : 0u); } /* - * r444: the entry is being consumed; no retry can need its - * staging snapshot any more (see svc_copy_snap_free_owner). - * - * r456: EXCEPT on the APPLY path. The apply runs at the head of a - * re-decode pass of an entry whose copy is still held for its - * operand pins; the pass's own op execution (which consumes the - * snapshot on success) follows the apply. Freeing here destroyed - * the pipeline-order snapshot before it could ever be read and - * the next pass's guard re-capture then read the driver's - * REFILLED ping-pong staging - the r455 operand-shift: the a1 H2D - * executed ~3 s after its push with a2's bytes (owner - * 0x5c00005c/0/6076682fa76b8792, snapshot first4 flipped from - * 00000000 3f000000... to 00000000 3e000000... across retries; - * dst16_after=000000000000003e0000803e0000c03e proves the a2 - * pattern landed in da1's page). The snapshot now survives the - * retries and is consumed (and freed) by the copy that finally - * executes; svc_copy_snap_free_owner still runs on every genuine - * consume path (FINISH/FLUSH/DROP/STALE/COUNTED/SWEEP/STOP) and - * the TTL covers the rest. + * r456/r457: the snapshot lifecycle is now "capture once, consume + * on the copy's success, evict on ring-slot reuse, TTL". Every + * opdef_clear caller can still re-decode the entry: + * - APPLY (0): the pass's own op execution (which consumes the + * snapshot on success) follows the apply; + * - FINISH (1): the entry can be REPLAYED after a consume (the + * pin-batch rewind replays entries while mappings stream in), + * and an owed copy still needs the pipeline-order bytes; + * - FLUSH (2): the op stays skipped this pass but a later replay + * may run it. + * The r457 mode-C evidence: the b1 H2D's snapshot was captured at + * 191607.816388 with b1's bytes (first4=00000000 3e800000 ...), + * FINISH freed it on the entry's consume, the rewind replay + * re-captured at 191608.798838 from the REFILLED staging + * (first4=00000000 3d800000 ... = b2) and the deferred execution + * wrote b2's bytes into db1 (dst16_after=000000000000803d...). + * The r444 leak the old free guarded against is already covered + * by the owner key (a stale snapshot can only ever serve its own + * (ch, idx, hash)), the slot-reuse eviction in + * svc_copy_snap_capture, the ops_done skip free in svc_run_ops + * and the 30 s TTL. */ - if (e->active && caller != SVC_OPDEF_CLR_APPLY) - svc_copy_snap_free_owner(e->h_object, e->entry_idx); memset(e, 0, sizeof(*e)); } @@ -21358,10 +21498,81 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) * through to the historical * consume so the channel can * never livelock. + * + * r470: an ALL-ZERO push page is + * the same "not the client's work + * yet" condition - the ring entry + * is written (the previous + * session's or the client's own + * slot reuse) but the push + * CONTENT has not landed, so the + * decode sees no QMD and the + * entry would be consumed as a + * no-op; the end-of-pass counter + * publish then advances the + * guest's completion word while + * the real push was never read, + * and the client drains a + * staging buffer that still holds + * the PREVIOUS copy's bytes + * (streams_events: h_c[2] = c0 + * because the c2 D2H's + * completion was published 4.5 s + * before its copy ran, the ring + * slots having been decoded as + * all-zero pushes). Deliberately + * NARROWER than the CE-shaped + * plausibility verdict: compute + * pushes legitimately fail that + * matcher, so only an entirely + * zero page counts as unwritten. */ + { + bool all_zero = false; + + if (!ch->engine_type && + va && len) { + u32 zi; + + all_zero = true; + for (zi = 0; zi < len; + zi++) + if (pb[zi]) { + all_zero = + false; + break; + } + } + /* + * r471: publish clamp bookkeeping. + * A zero-content GR entry that is + * held or consumed must never let + * the end-of-pass counter publish + * advance the guest's completion + * word past it (the client would + * satisfy its readback wait before + * the real push/copy ran - the + * staging buffer still holds the + * previous copy's bytes). Record + * the lowest zero slot; a decode + * that reads non-zero content + * clears the clamp. + */ + if (!ch->engine_type) { + if (all_zero) { + if (!ch->gr_bad_put || + (u32)idx + 1 < + ch->gr_bad_put) + ch->gr_bad_put = + (u32)idx + 1; + } else if (va && len) { + ch->gr_bad_put = 0; + } + } if (!ch->engine_type && va && len && gsp_phys_gr_forward_enabled() && - svc_gr_qmd_seen > 0 && + (svc_gr_qmd_seen > 0 || + all_zero) && !entry_replay && !(svc_gr_cand.valid && svc_gr_cand.h_object == @@ -21413,6 +21624,7 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) ch->h_object, (u32)idx, va, len); } + } /* * A CE push that should carry this * channel's next tracking release but @@ -22148,6 +22360,10 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) svc.chan[k].push_hold_until = ch->push_hold_until; svc.chan[k].defer_since = ch->defer_since; svc.chan[k].put_prev = ch->put_prev; + /* r471: the publish clamp watermark is set + * and cleared on the pump thread; persist it + * so the next pass still clamps. */ + svc.chan[k].gr_bad_put = ch->gr_bad_put; /* * Persist the cheap-wait window and the content * watermark: the pump works on a snapshot, so 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/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/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..5fcdec0 --- /dev/null +++ b/tools/cuda-ab/run/tgemm_probe 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/guest-diag/pin_gr_stream.py b/tools/guest-diag/pin_gr_stream.py index 6ba51ee..6eb5be7 100644 --- a/tools/guest-diag/pin_gr_stream.py +++ b/tools/guest-diag/pin_gr_stream.py @@ -211,7 +211,7 @@ def _demand_loop(): except Exception: continue now = time.time() - for m in re.finditer(r"(?:stalled VA #\d+|pending release starved va)=(0x[0-9a-f]+)", + 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]) |