diff options
34 files changed, 3724 insertions, 101 deletions
diff --git a/docs/03-architecture/13-vexecution-plane.md b/docs/03-architecture/13-vexecution-plane.md index caaad01..3d97d07 100644 --- a/docs/03-architecture/13-vexecution-plane.md +++ b/docs/03-architecture/13-vexecution-plane.md @@ -568,6 +568,39 @@ forwards (local evidence and the prepared harness: forwarding-service cost, not engine time — the number to drive down is the 7.9 ms and the per-window round trips. +### 3.5a Copy-out change detection: fingerprint rule + +A forwarded launch stages every referenced data page into the data-bounce +chunk, and after the retire wait the copy-out ships back only the pages the +engine actually modified. That decision used to be a 32-bit additive +checksum of the page's words, compared against a baseline captured at +copy-in. The additive form has a structural blind spot: doubling every word +of a 4 KB page raises each float exponent by one (+2^23 per word), and 1024 +words contribute `1024 * 2^23 = 2^33`, which is exactly `0 (mod 2^32)` — an +in-place power-of-two scale of a whole page (`c[i] *= 2.0f`) left the +checksum bit-identical, the page was classified "engine did not touch", +the write-back was skipped, and the client read the pre-scale values. The +same blindness applies to `*= 0.5f`, `* 4`, `* 8`, etc. Since `r452` +(commit `19399f4f`) every change decision in the copy path — the retire +wait's settle scan, the tracked and gap copy-outs, the window witness, the +late-store witness and the idempotent-launch fingerprint — uses a 64-bit +FNV-1a hash over the page instead; the additive sum survives only as the +all-zero filter. Policy rule for the plane: **a change-detection +fingerprint must not be an algebraic invariant of any plausible kernel +operation** — use a hash, never a sum. + +The reader side of the same decision is the landing table: a D2H copy of a +device VA resolves through the last recorded write of that VA. Since `r453` +the forwarded QMD's release0 is redirected to the module semaphore by +default (`gsp_phys_rmcall_gr_rel_mod=1`): the guest-visible release word is +authored by the host only on the completion edge, after the copy-out has +shipped the engine's stores, so a client's stream-ordered readback can no +longer run ahead of the write-back (the historical shape: the engine's own +release satisfied the guest while the copy-out was still pending). The +engine-completion observable remains the module semaphore, and the +per-outcome authoring path guarantees the guest word is applied on every +outcome (timeout, wedge, no-write). + ### 3.6 Citation map for this section | Claim | Authority | diff --git a/docs/05-operations/06-guest-frame-path-and-looking-glass.md b/docs/05-operations/06-guest-frame-path-and-looking-glass.md index 3dda4f1..bc5447a 100644 --- a/docs/05-operations/06-guest-frame-path-and-looking-glass.md +++ b/docs/05-operations/06-guest-frame-path-and-looking-glass.md @@ -81,3 +81,22 @@ python3 tools/lg/frame_reader.py --frames 5 --out /tmp/opencode/frames gcc -O2 -o egl_frame tools/guest-diag/egl_frame.c -lEGL -lGLESv2 ./egl_frame ``` + +## 5. Session-13 update (2026-10-21, r432→r441) + +* The single-frame path is unchanged and still proven: guest `frame_producer` + (CUDA Julia) -> kvmfr ivshmem BAR -> host `frame_reader.py` -> + `/tmp/opencode/frames/frame-0001.png` (320x240, 1431 colours). +* `tools/lg/frame_sequence.sh` captures an **animated** sequence end to end + (reader first, then N frames from the guest). It is blocked by the + GR no-execute wedge: when the lane is wedged the render launches write zero + pages and the producer republishes stale identical frames (observed: 6/6 + identical MD5s on r438; r439 published 5 identical plus one frame left in + the BAR by the previous boot). A multi-frame animation needs a healthy + lane; the frame pipeline itself is not the blocker. +* The GL wall is untouched: `Xorg :0` -> `(EE) NVIDIA(GPU-0): Failed to + allocate VASpace object`, guest dmesg `NVRM: vaspaceapiConstruct_IMPL: + Could not construct VA space. Status ffff`, i.e. `vmmCreateVaspace` inside + the guest's own RM. Next step: reproduce with the guest RM's + `ResmanDebugLevel` raised (a guest module reload, the service's m177 path) + to capture which inner call returns the generic error. diff --git a/kernel/gsp_phys_booter.c b/kernel/gsp_phys_booter.c index 5350f90..42a0f2a 100644 --- a/kernel/gsp_phys_booter.c +++ b/kernel/gsp_phys_booter.c @@ -140,6 +140,7 @@ static inline bool gsp_phys_gr_copywit_enabled(void) return READ_ONCE(gsp_phys_gr_copyout_dump_param); } + /* * r405 late-write probes (the r405 "partial write" class: the scale launch * of streams2 wrote 1 of its 32 expected in-place pages and every other @@ -9626,41 +9627,66 @@ static int gsp_phys_rmcall_golden_replay(void __iomem *bar0) * (swref/published/pascal/gp100/dev_mmu.h:97-119,157; * uvm_pascal_mmu.c:64-141,235-303). */ -static u64 gsp_phys_rmcall_pde(u64 pa) -{ - return 0xeull | ((pa >> 12) << 8); -} - /* - * r431 COHERENCY ROOT-CAUSE CANDIDATE: the leaf-PTE aperture for every host - * page this backend maps (the ladder windows, the win0 shadow pool, the + * r431/r436 COHERENCY ROOT-CAUSE CANDIDATE: the leaf-PTE aperture for every + * host page this backend maps (the ladder windows, the win0 shadow pool, the * bounce/data regions, the release/scratch shadows) was hardcoded to - * aperture 3 = SYSTEM_NON_COHERENT_MEMORY. The reference RM picks the - * aperture from the allocation's CPU cache attribute - * (fabric_vaspace.c:1123-1134): NV_MEMORY_CACHED -> SYSTEM_COHERENT (2), - * otherwise NON_COHERENT (3). Every page here is dma_alloc_coherent (CPU - * cached/coherent by construction), so 3 is wrong: the GPU's L2 does not - * snoop the CPU, which makes the engine's stores linger in L2 (the host's - * retire scan reads stale memory = the "no-execute wedge") and the host's - * re-staged copy-in content invisible to the engine (the retries re-read - * stale lines). gsp_phys_rmcall_pte_coh selects aperture 2 (default 1) for - * a one-flag A/B against the historical 3 (0); the PTEs are planted before - * the engine's first fetch, so the flag is boot-time only. - */ -static bool gsp_phys_rmcall_pte_coh_param = true; + * aperture 3 = SYSTEM_NON_COHERENT_MEMORY. The reference RM derives the + * pair from the allocation's CPU cache attribute (fabric_vaspace.c:1123-1134 + * for the aperture, mem_desc.c memdescGetVolatility for VOL: NV_MEMORY_CACHED + * -> SYSTEM_COHERENT (2) + volatile FALSE, NV_MEMORY_UNCACHED -> NON_COHERENT + * (3) + volatile TRUE). The r431 experiment used aperture 2 with VOL still 1 + * (0xd) - a pair the reference RM never produces - and it made the engine's + * truncated writes visible sooner but worsened the wedge on a fresh boot. + * r436 turns the knob into the full 2x2 matrix so each consistent encoding + * can be A/B'd: + * 0 = aperture 3 + VOL 1 (0xf) - historical, NV_MEMORY_UNCACHED + * 1 = aperture 2 + VOL 1 (0xd) - r431 experiment (inconsistent pair) + * 2 = aperture 2 + VOL 0 (0x5) - NV_MEMORY_CACHED (coherent, non-volatile) + * 3 = aperture 3 + VOL 0 (0x7) - non-coherent, non-volatile + * The PTEs are planted before the engine's first fetch, so the flag is + * boot-time only. + */ +static int gsp_phys_rmcall_pte_coh_param; module_param_named(gsp_phys_rmcall_pte_coh, gsp_phys_rmcall_pte_coh_param, - bool, 0444); + int, 0444); MODULE_PARM_DESC(gsp_phys_rmcall_pte_coh, - "r431: sysmem leaf-PTE aperture: 1 = SYSTEM_COHERENT (correct for dma_alloc_coherent host pages), 0 = SYSTEM_NON_COHERENT (historical)"); + "r436: sysmem leaf-PTE aperture/VOL mode: 0=0xf (aperture3+VOL, historical), 1=0xd (aperture2+VOL, r431), 2=0x5 (aperture2, non-volatile), 3=0x7 (aperture3, non-volatile)"); + +static u64 gsp_phys_rmcall_sysmem_base_low(void) +{ + switch (gsp_phys_rmcall_pte_coh_param) { + case 1: + return 0xdull; /* aperture 2, VOL 1 */ + case 2: + return 0x5ull; /* aperture 2, VOL 0 */ + case 3: + return 0x7ull; /* aperture 3, VOL 0 */ + default: + return 0xfull; /* aperture 3, VOL 1 (historical) */ + } +} static u64 gsp_phys_rmcall_pte(u64 pa) { /* * Low nibble: VALID (1) | APERTURE<<1 | VOL (8). Aperture 2 = * SYSTEM_COHERENT_MEMORY -> 1|4|8 = 0xd; aperture 3 = NON_COHERENT - * -> 1|6|8 = 0xf. + * -> 1|6|8 = 0xf. r436: see gsp_phys_rmcall_pte_coh_param. */ - u64 base = gsp_phys_rmcall_pte_coh_param ? 0xdull : 0xfull; + u64 base = gsp_phys_rmcall_sysmem_base_low(); + + return base | ((pa >> 12) << 8); +} + +/* + * r436: the single PDE uses the same aperture/VOL pair as the leaf PTEs but + * with IS_PDE=0 in bit 0 (the low bit of the PTE base is VALID there), so the + * PDE base is the PTE base with bit 0 cleared: 0xe/0xc/0x4/0x6. + */ +static u64 gsp_phys_rmcall_pde(u64 pa) +{ + u64 base = gsp_phys_rmcall_sysmem_base_low() & ~1ull; return base | ((pa >> 12) << 8); } @@ -15123,6 +15149,39 @@ static int gsp_phys_rmcall_ce_submit_locked(void __iomem *bar0, u64 src_va, * semaphore. On timeout the existing status-queue drain runs and the helper * returns -ETIMEDOUT; it never retries. */ +static int gsp_phys_rmcall_gr_resched_locked(void __iomem *bar0); +static int gsp_phys_rmcall_gr_sem_push_locked(void __iomem *bar0); + +/* + * r445 wedge hard re-arm: the outcome-retry's warm semaphore probe proves the + * engine refused the channel's proven GR block (the sticky post-dispatch + * no-store wedge). One bounded recovery attempt instead of an immediate + * break: force a context save/unload with a bWait TSG preempt, re-run the + * mode-35 lane-enable tail + schedules (the m155 extended re-arm), then + * repeat the warm probe. Caller holds gsp_phys_rm_lock. Returns 0 when the + * engine executed the probe after the re-arm. + */ +static int gsp_phys_gr_wedge_hard_rearm_locked(void __iomem *bar0) +{ + u32 pre[2] = { 1u, 0u }; /* bWait=1, bManualTimeout=0 */ + u32 rpc_result = 0, status = 0; + int ret; + + ret = gsp_phys_rmcall_control(bar0, boot_stage.rmcall_tsg, + GSP_PHYS_RMCALL_CTRL_TSG_PREEMPT, + pre, sizeof(pre), sizeof(pre), NULL, 0, + &rpc_result, &status, + "TSG PREEMPT bWait=1 (r445 wedge hard re-arm)"); + pr_warn("vgsp-gr: WEDGE-REARM preempt status=0x%x rpc=0x%x ret=%d\n", + status, rpc_result, ret); + ret = gsp_phys_rmcall_gr_resched_locked(bar0); + pr_warn("vgsp-gr: WEDGE-REARM resched ret=%d\n", ret); + ret = gsp_phys_rmcall_gr_sem_push_locked(bar0); + pr_warn("vgsp-gr: WEDGE-REARM probe rc=%d (engine %s after re-arm)\n", + ret, ret ? "still wedged" : "ALIVE"); + return ret; +} + static int gsp_phys_rmcall_gr_sem_push_locked(void __iomem *bar0) { struct gsp_phys_boot_chunk *userd_c = @@ -16309,6 +16368,19 @@ EXPORT_SYMBOL_GPL(gsp_phys_backend_guest_ce_copy); #define GSP_PHYS_RMCALL_GR_CWD_REFCOUNT 1536u #define GSP_PHYS_RMCALL_GR_CWD_LEVELS 64u #define GSP_PHYS_RMCALL_GR_MAX_SM_COUNT 24u +/* + * r434 CWD-slot-count path: the falcon PRI write (CUDA's + * turingFalconPRIWriteMasked: flag/data/mask through SET_MME_SHADOW_SCRATCH + * 0/1/2, then SET_FALCON04 = register offset) and the gated method. Used only + * when gsp_phys_rmcall_gr_cwd_slot >= 0. + */ +#define GSP_PHYS_RMCALL_GR_MME_SCRATCH0 0x3400u +#define GSP_PHYS_RMCALL_GR_MME_SCRATCH1 0x3404u +#define GSP_PHYS_RMCALL_GR_MME_SCRATCH2 0x3408u +#define GSP_PHYS_RMCALL_GR_SET_FALCON04 0x0510u +#define GSP_PHYS_RMCALL_GR_SET_CWD_SLOT_COUNT 0x02b0u +#define GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1 0x00407040u +#define GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1_EN_CWD_SLOT_COUNT 0x800u /* Method block bound: init(152) + inline run(67) + PCAS(4) + REPORT_SEM(8) + * the optional QMD-virt pair(2). */ #define GSP_PHYS_RMCALL_GR_METHOD_WORDS 512u @@ -16401,7 +16473,7 @@ static bool gsp_phys_rmcall_gr_init_emitted; * stale-QMD execution) landed after its launch was shipped. */ static u32 gsp_phys_gr_prev_out_bi; -static u32 gsp_phys_gr_prev_out_csum; +static u64 gsp_phys_gr_prev_out_csum; static bool gsp_phys_gr_prev_out_valid; /* r425: SKED/FECS/CWD register witness (one line; the healthy-vs-wedged @@ -16529,6 +16601,29 @@ static u32 *gsp_phys_rmcall_build_gr_init(u32 *p) GSP_PHYS_RMCALL_GR_MAX_SM_COUNT); p = gsp_phys_rmcall_push_gr(p, GSP_PHYS_RMCALL_GR_INIT_SM_DISP_CTRL, 0u); + /* + * r434 CWD-slot-count path (gsp_phys_rmcall_gr_cwd_slot >= 0): masked + * PRI write of SKED_DEBUG_1.ENABLE_CWD_SLOT_COUNT (bit 11) through the + * falcon scratch registers, then the gated SET_CWD_SLOT_COUNT method - + * the CUDA driver's own turingLaunchSetCwdSlotCount sequence. -1 (the + * default) emits nothing, so the historical preamble is unchanged. + */ + if (gsp_phys_rmcall_gr_cwd_slot_value() >= 0) { + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_MME_SCRATCH0, 0u); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_MME_SCRATCH1, + GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1_EN_CWD_SLOT_COUNT); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_MME_SCRATCH2, + GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1_EN_CWD_SLOT_COUNT); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_SET_FALCON04, + GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_SET_CWD_SLOT_COUNT, + (u32)gsp_phys_rmcall_gr_cwd_slot_value()); + } return p; } @@ -17295,6 +17390,54 @@ static struct gsp_phys_boot_chunk *gsp_phys_gr_data_chunk(void) } /* + * r452: scan the whole active GR data-bounce chunk for an 8-byte pattern + * (write "0x<pattern>" to /sys/module/vgsp/parameters/gsp_phys_gr_scan). + * Verifies by content where a lost engine store actually landed: the chunk + * is the only bulk memory the operand window PTEs map, so a store that + * executed but is absent from the writer's tracked pages must be findable + * here (a stale-PTE landing) or nowhere (a dropped store). + */ +static int gsp_phys_gr_scan_set(const char *val, + const struct kernel_param *kp) +{ + struct gsp_phys_boot_chunk *c = gsp_phys_gr_data_chunk(); + u64 pattern = 0; + u32 i, hits = 0; + + if (kstrtoull(val, 0, &pattern)) + return -EINVAL; + if (!c->pages) + return -ENODEV; + for (i = 0; i < c->pages; i++) { + const u8 *p = gsp_phys_chunk_page_va(c, i); + u32 off; + + if (!p) + continue; + for (off = 0; off + 8 <= GSP_PHYS_PAGE_SIZE; off += 8) { + if (get_unaligned_le64(p + off) == pattern) { + if (hits < 32) + pr_info("vgsp-gr: CHUNK-SCAN found 0x%llx at bi=%u off=0x%x (va=0x%llx)\n", + pattern, i, off, + GSP_PHYS_RMCALL_GR_DATA_BOUNCE_VA0 + + ((u64)i << 12) + off); + hits++; + } + } + } + pr_info("vgsp-gr: CHUNK-SCAN 0x%llx done: %u hit(s) over %u pages\n", + pattern, hits, c->pages); + return 0; +} + +static const struct kernel_param_ops gsp_phys_gr_scan_ops = { + .set = gsp_phys_gr_scan_set, +}; +module_param_cb(gsp_phys_gr_scan, &gsp_phys_gr_scan_ops, NULL, 0200); +MODULE_PARM_DESC(gsp_phys_gr_scan, + "Scan the GR data-bounce chunk for an 8-byte pattern (hex)"); + +/* * B-2: page index of a relocated operand VA inside the engine-visible data * region, replacing the five inlined (reloc_va - 0x110000) >> 12 derivations * and their MAIL_SLOT bounds. Grown region: VAs from @@ -17744,7 +17887,7 @@ static bool gsp_phys_gr_reloc_fix_word(const struct gsp_phys_gr_submit *req, */ static noinline u32 gsp_phys_gr_scan_pages(const struct gsp_phys_gr_submit *req, - const u32 *csum_in, u32 *n_tracked, u64 *hash_out) + const u64 *csum_in, u32 *n_tracked, u64 *hash_out) { static u32 checked[GSP_PHYS_RMCALL_GR_DATA_BOUNCE_PAGES]; struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); @@ -17790,15 +17933,24 @@ gsp_phys_gr_scan_pages(const struct gsp_phys_gr_submit *req, *n_tracked = n_all; for (j = 0; j < n_all; j++) { const u8 *src = gsp_phys_chunk_page_va(dchunk, checked[j]); - u32 ci, csum = 0; + u64 fp; if (!src) continue; - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(src + ci); - if (csum != csum_in[checked[j]]) + /* + * r452: FNV-1a per page, not a 32-bit additive sum. The + * add-sum was PROVABLY blind to a power-of-two in-place + * scale: doubling every word of a 4 KB page shifts each + * exponent by +1, i.e. +2^23 per word, and 1024 words give + * 2^33 == 0 (mod 2^32) - so c[i] *= 2.0f left the checksum + * bit-identical, the page classified "unchanged", and the + * client read the pre-scale values (the deterministic + * streams2/scale class). FNV-1a has no such structure. + */ + fp = gsp_phys_gr_fnv1a64(src, GSP_PHYS_PAGE_SIZE); + if (fp != csum_in[checked[j]]) changed++; - hash ^= csum; + hash ^= fp; hash *= 1099511628211ULL; } if (hash_out) @@ -17806,9 +17958,35 @@ gsp_phys_gr_scan_pages(const struct gsp_phys_gr_submit *req, return changed; } +/* + * r434: post-launch deposit dump (hardware-only words). The r430 QMD64 dump + * runs at submission, before the engine touches the deposit, so the + * hardware-only fields (w50 INNER_GET, w51 INNER_PUT, w52 span-list head, + * w53 next-QMD) are always zero there. Reading the same page AFTER the + * retire wait shows what the SKED actually consumed: w50/w51 are the CTA + * (INNER) get/put counters, so a wedged launch whose grid "completed" with + * zero stores is distinguishable from a genuinely dispatched grid by + * whether INNER_PUT advanced to the CTA count. One line per launch, gated + * by the skedrec param. + */ +static void gsp_phys_gr_qmd64_post(const u8 *deposit, const char *tag) +{ + char qb[64 * 9 + 1]; + int qo = 0; + u32 i; + + if (!deposit || !gsp_phys_rmcall_gr_skedrec_value()) + return; + for (i = 0; i < GSP_PHYS_GR_QMD_WORDS && qo < (int)sizeof(qb) - 10; + i++) + qo += scnprintf(qb + qo, sizeof(qb) - qo, "%08x ", + get_unaligned_le32(deposit + i * 4)); + pr_info("vgsp-gr: QMD64-POST %s %s\n", tag, qb); +} + static noinline u32 gsp_phys_gr_output_retire_wait(const struct gsp_phys_gr_submit *req, - const u32 *csum_in, u32 *n_tracked) + const u64 *csum_in, u32 *n_tracked) { u32 n_all = 0, iter; u32 changed = 0; @@ -17962,7 +18140,7 @@ static u32 gsp_phys_gr_idem_n; static u64 gsp_phys_gr_launch_fp(const u32 *qmd, const struct gsp_phys_gr_submit *req, - const u32 *csum_in) + const u64 *csum_in) { u64 h = 1469598103934665603ULL; u32 i; @@ -18040,6 +18218,22 @@ static void gsp_phys_gr_page_spot_check(const u8 *dst, const u8 *src, } /* + * r452: the legacy 32-bit additive page sum, kept ONLY for the all-zero + * filter. It is never a change-detection fingerprint any more: doubling every + * word of a 4 KB page leaves the sum bit-identical (1024 x 2^23 == 2^33 == 0 + * mod 2^32), which silently classified an in-place x2 scale as "no engine + * store" and shipped stale values (the deterministic streams2 class). + */ +static inline u32 gsp_phys_gr_page_sum32(const u8 *p) +{ + u32 ci, csum = 0; + + for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) + csum += get_unaligned_le32(p + ci); + return csum; +} + +/* * COPYWIT r226 zero filter: an all-zero page is never a meaningful alias * partner. A fresh client's fill stages long runs of zero pages * (uninitialized cudaMalloc buffers), and consecutive zero pages produced @@ -18085,7 +18279,7 @@ static u32 gsp_phys_gr_in_census_log; * again). Returns the number of pages staged. */ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, - u32 *csum_in, u64 *staged_hpa) + u64 *csum_in, u64 *staged_hpa) { struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); struct { @@ -18122,7 +18316,9 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, "gr-bounce-copy-in"); for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) csum += get_unaligned_le32(dst + ci); - csum_in[bi] = csum; + /* r452: the change baseline is the FNV-1a fingerprint (the + * add-sum above survives only as the all-zero test). */ + csum_in[bi] = gsp_phys_gr_fnv1a64(dst, GSP_PHYS_PAGE_SIZE); zero = gsp_phys_gr_page_all_zero(csum, dst); if (!zero && gsp_phys_gr_copywit_enabled()) { u32 j; @@ -18171,7 +18367,7 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, * retry, so a retried launch starts from the guest inputs in both sets. * Returns the number of pages staged. */ -static u32 gsp_phys_gr_fill_gap_bounce(u32 *csum_in, u64 *staged_hpa) +static u32 gsp_phys_gr_fill_gap_bounce(u64 *csum_in, u64 *staged_hpa) { struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); u32 gi, copied = 0; @@ -18179,7 +18375,6 @@ static u32 gsp_phys_gr_fill_gap_bounce(u32 *csum_in, u64 *staged_hpa) for (gi = 0; gi < gsp_phys_gr_gap_n; gi++) { const u8 *src = (const u8 *)__va(gsp_phys_gr_gap[gi].hpa); u8 *dst = gsp_phys_chunk_page_va(dchunk, gsp_phys_gr_gap[gi].bi); - u32 ci, csum = 0; if (!src || !dst) continue; @@ -18189,9 +18384,9 @@ static u32 gsp_phys_gr_fill_gap_bounce(u32 *csum_in, u64 *staged_hpa) staged_hpa[gsp_phys_gr_gap[gi].bi] = gsp_phys_gr_gap[gi].hpa; gsp_phys_gr_page_spot_check(dst, src, gsp_phys_gr_gap[gi].va, "gr-bounce-copy-in"); - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(dst + ci); - csum_in[gsp_phys_gr_gap[gi].bi] = csum; + /* r452: FNV-1a baseline (see gsp_phys_gr_scan_pages). */ + csum_in[gsp_phys_gr_gap[gi].bi] = + gsp_phys_gr_fnv1a64(dst, GSP_PHYS_PAGE_SIZE); copied++; } dma_wmb(); @@ -18454,6 +18649,8 @@ static u32 gsp_phys_m91_fault_start; static u32 gsp_phys_m91_cap_used; static bool gsp_phys_m91_cap_suppressed; static bool gsp_phys_m91_icache_log; +static bool gsp_phys_r449_ccache_log; +static bool gsp_phys_gr_rel_membar_log; static noinline void gsp_phys_gr_m91_fault_cap(const struct gsp_phys_gr_submit *req, const u32 *qmd, u64 qmd_dst, const u8 *qmd_dst_mem, @@ -18785,7 +18982,7 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, * stack). Safe as static: the only caller is the single-inflight * ordered gsp-phys-work item. */ - static u32 bounce_csum_in[GSP_PHYS_RMCALL_GR_DATA_BOUNCE_PAGES]; + static u64 bounce_csum_in[GSP_PHYS_RMCALL_GR_DATA_BOUNCE_PAGES]; /* * r378: which guest page's content each bounce slot holds this * launch (set by the copy-in fills, cleared per submission). The @@ -18963,6 +19160,52 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, } } /* + * r449 SHADER-CONSTANT-CACHE invalidate (gsp_phys_gr_ccache_inv; default + * 0 = off): the same cache-control word's bit 31 (the V01_06-layout + * SHADER_CONSTANT_CACHE position) is never set by the captured guest + * QMDs. The forwarded param page is rewritten in place between launches + * (the two-bank win0 pool), so a stale constant-cache line can serve + * the PREVIOUS launch's parameters and the current launch's output + * pages stay untouched - the post-dispatch no-store signature. Mode 1 + * sets bit 31; mode 2 forces the whole invalidate set (bits 24..31). + * Only when the QMD carries the driver's standard cache-control + * signature (bits 26..29) so an unknown QMD revision is never touched. + */ + if (gsp_phys_gr_ccache_inv_value() && + (qmd[5] & 0x3c000000u) == 0x3c000000u) { + u32 mode = gsp_phys_gr_ccache_inv_value(); + + if (mode == 2) + qmd[5] |= 0xff000000u; + else + qmd[5] |= 1u << 31; + if (!gsp_phys_r449_ccache_log) { + gsp_phys_r449_ccache_log = true; + pr_info("vgsp-gr: r449 constant-cache invalidate mode=%u applied (w5=0x%08x)\n", + mode, qmd[5]); + } + } + /* + * r433 RELEASE_MEMBAR_TYPE A/B (gsp_phys_gr_rel_membar): force the + * forwarded QMD's RELEASE_MEMBAR_TYPE to FE_SYSMEMBAR (QMDV00_06 + * word 11 bit 14) so the front-end flushes its pending system-memory + * stores before writing the release semaphore. The captured guest + * QMDs carry FE_NONE, so a grid that completes with its data still + * in the GPU's L2 releases the client while no store is host-visible + * - the post-dispatch store-loss signature. The release ordering then + * guarantees the host's retire scan and copy-out can never observe a + * completion whose data has not landed. 0 (default) keeps the + * guest's captured value. + */ + if (gsp_phys_gr_rel_membar_value() && !(qmd[11] & (1u << 14))) { + qmd[11] |= 1u << 14; + if (!gsp_phys_gr_rel_membar_log) { + gsp_phys_gr_rel_membar_log = true; + pr_info("vgsp-gr: r433 RELEASE_MEMBAR_TYPE forced to FE_SYSMEMBAR in the forwarded QMD (w11=0x%08x)\n", + qmd[11]); + } + } + /* * Two-step PCAS (gsp_phys_rmcall_gr_pcas_two_step): set the QMD's * REQUIRE_SCHEDULING_PCAS (w4 bit12) so the copy PCAS does not * schedule and the trailing schedule-only PCAS does. @@ -19789,6 +20032,27 @@ gr_page_unmapped: /* A stale release from a previous submission must not be observed. */ WRITE_ONCE(*(u32 *)sem_va, 0); WRITE_ONCE(*(u32 *)(sem_va + 4), 0); + /* + * r441: zero the GUEST-VISIBLE release words for the first attempt + * too, not only for outcome retries. The r439 evidence shows a first + * attempt whose poll broke on a STALE value (qmd_rel=0x00000005 from + * the previous launch's engine write, want=6) while the engine never + * fetched this launch's deposit (sentinel intact, "engine fetch + * NONE"): the host then shipped a completion for a launch that did + * not run. Zeroing forces the poll to wait for THIS launch's own + * engine write; when the engine does not write it, the timeout ladder + * and the software-authored completion take over (the honest path), + * and the wedge classification stops being satisfiable by a stale + * value. The write guard re-validates the target like the retry path. + */ + if (qmd_rel_hpa && qmd_rel_off < GSP_PHYS_PAGE_SIZE && + gsp_sema_gr_write_guard(req->rel_va, qmd_rel_hpa, + "gr-rel-zero-first")) + WRITE_ONCE(*(u32 *)((u8 *)__va(qmd_rel_hpa) + qmd_rel_off), 0); + if (ch_rel_hpa && ch_rel_off < GSP_PHYS_PAGE_SIZE && + gsp_sema_gr_write_guard(req->rel_va, ch_rel_hpa, + "gr-chrel-zero-first")) + WRITE_ONCE(*(u32 *)((u8 *)__va(ch_rel_hpa) + ch_rel_off), 0); wmb(); /* @@ -20717,6 +20981,14 @@ gr_page_unmapped: &n_tracked); /* + * r434: post-launch hardware-only QMD dump (w50/w51 INNER + * get/put = the SKED's CTA dispatch counters) so the + * healthy-vs-wedged diff has the fields the submission-time + * dump cannot show. + */ + gsp_phys_gr_qmd64_post(qmd_dst_mem, "retire"); + + /* * m179: the zero-write test is specific to launches with a * non-empty TRACKED data set. A launch whose data entries do * not resolve to any bounce page legitimately reports zero @@ -20805,6 +21077,26 @@ gr_page_unmapped: pr_info("vgsp-gr: outcome retry %u/%u (wrote=%u tracked=%u): refill + re-publish\n", nw_attempts, nw_retries, wrote, n_tracked); + /* + * r445 wedge hard re-arm (gsp_phys_gr_wedge_rearm, + * default off): the r445 evidence shows every + * post-dispatch no-store episode leaves the warm + * ladder sem probe ALIVE (the ladder block executes) + * while the client launch keeps writing nothing, so + * the plain refill+re-publish retries are futile. + * On the FINAL retry attempt, force one context + * save/unload (TSG PREEMPT bWait=1) and re-run the + * mode-35 lane-enable tail + schedules before the + * re-publish, so the retried launch meets a freshly + * restored context. Bounded: one preempt per launch + * (2 s RPC timeout), only on the last retry. + */ + if (nw_attempts == nw_retries && + gsp_phys_gr_wedge_rearm_value()) { + mutex_lock(&gsp_phys_rm_lock); + gsp_phys_gr_wedge_hard_rearm_locked(bar0); + mutex_unlock(&gsp_phys_rm_lock); + } if (quiet_ms) { gsp_phys_gr_wedge_quiet++; pr_warn("vgsp-gr: WEDGE-QUIET hold %u ms before outcome retry %u (push=0x%llx; %u consecutive zero-write launch(es); no submissions in between)\n", @@ -21085,6 +21377,8 @@ gr_page_unmapped: } wrote = gsp_phys_gr_output_retire_wait( req, bounce_csum_in, &n_tracked); + gsp_phys_gr_qmd64_post(qmd_dst_mem, + "retry-retire"); pr_info("vgsp-gr: outcome retry %u done: engine wrote %u page(s)\n", nw_attempts, wrote); /* r410: the quiet-window verdict. */ @@ -21222,14 +21516,12 @@ gr_page_unmapped: const u8 *pp = gsp_phys_chunk_page_va( gsp_phys_gr_data_chunk(), gsp_phys_gr_prev_out_bi); - u32 pc = 0, pi; + u64 pc; if (pp) { - for (pi = 0; pi < GSP_PHYS_PAGE_SIZE; - pi += 4) - pc += get_unaligned_le32(pp + pi); + pc = gsp_phys_gr_fnv1a64(pp, GSP_PHYS_PAGE_SIZE); if (pc != gsp_phys_gr_prev_out_csum) - pr_warn("vgsp-gr: PREV-OUT-CHANGED bi=%u csum %08x -> %08x (late store after its launch shipped)\n", + pr_warn("vgsp-gr: PREV-OUT-CHANGED bi=%u fp %016llx -> %016llx (late store after its launch shipped)\n", gsp_phys_gr_prev_out_bi, gsp_phys_gr_prev_out_csum, pc); @@ -21326,7 +21618,7 @@ gr_page_unmapped: for (wbi = 0; wbi < req->n_data; wbi++) { u32 bi = gsp_phys_gr_data_bi(req->data[wbi].reloc_va); const u8 *src; - u32 ci, csum = 0; + u64 fp; if (bi == GSP_PHYS_GR_DATA_BI_NONE || !req->data[wbi].hpa) @@ -21347,9 +21639,8 @@ gr_page_unmapped: aliased++; continue; } - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(src + ci); - if (csum == bounce_csum_in[bi]) { + fp = gsp_phys_gr_fnv1a64(src, GSP_PHYS_PAGE_SIZE); + if (fp == bounce_csum_in[bi]) { /* * r378: the engine did not touch this page, * so the guest page still holds exactly this @@ -21358,6 +21649,37 @@ gr_page_unmapped: * stale-content corruption at worst - skip. */ unchanged++; + /* + * r452: witness the UNCHANGED page's chunk + * content (first 4 words) and the guest + * page's, bounded per launch. This is the + * only place the two can be compared for a + * page the csum rule classified as "engine + * did not touch" - it separates "the store + * never landed in the chunk" from "the + * store landed but the csum rule missed + * it" (or "the reader resolves a different + * page than the writer wrote"). + */ + if (gsp_phys_gr_copywit_enabled() && + outwit < GSP_PHYS_GR_COPYWIT_MAX) { + const u8 *gp = req->data[wbi].hpa ? + (const u8 *)__va(req->data[wbi].hpa) : + NULL; + + outwit++; + pr_info("vgsp-gr: OUT-UNCH bi=%u va=0x%llx fp=%016llx pre=%016llx chunk=%08x %08x %08x %08x guest=%08x %08x %08x %08x\n", + bi, req->data[wbi].va, fp, + bounce_csum_in[bi], + get_unaligned_le32(src), + get_unaligned_le32(src + 4), + get_unaligned_le32(src + 8), + get_unaligned_le32(src + 12), + gp ? get_unaligned_le32(gp) : 0, + gp ? get_unaligned_le32(gp + 4) : 0, + gp ? get_unaligned_le32(gp + 8) : 0, + gp ? get_unaligned_le32(gp + 12) : 0); + } continue; } { @@ -21408,13 +21730,13 @@ gr_page_unmapped: &rap, &rval); outwit++; - pr_info("vgsp-gr: OUT-PAGE-WIT va=0x%llx hpa=0x%llx bi=%u first8=%016llx last8=%016llx pre=%08x post=%08x res=%s pa=0x%llx ap=%u val8=%016llx vbase=%px\n", + pr_info("vgsp-gr: OUT-PAGE-WIT va=0x%llx hpa=0x%llx bi=%u first8=%016llx last8=%016llx pre=%016llx post=%016llx res=%s pa=0x%llx ap=%u val8=%016llx vbase=%px\n", req->data[wbi].va, req->data[wbi].hpa, bi, f8, (u64)get_unaligned_le64( src + GSP_PHYS_PAGE_SIZE - 8), - bounce_csum_in[bi], csum, + bounce_csum_in[bi], fp, rhave ? "ledger" : "none", rpa, rap, rval, (void *)gsp_sema_gr_shadow_base()); @@ -21429,7 +21751,8 @@ gr_page_unmapped: * evidence). */ if (gsp_phys_gr_copywit_enabled() && - !gsp_phys_gr_page_all_zero(csum, src)) { + !gsp_phys_gr_page_all_zero( + gsp_phys_gr_page_sum32(src), src)) { u32 wi; for (wi = 0; wi < gsp_phys_gr_outwit_n; @@ -21465,7 +21788,7 @@ gr_page_unmapped: */ if (changed == 0) { gsp_phys_gr_prev_out_bi = bi; - gsp_phys_gr_prev_out_csum = csum; + gsp_phys_gr_prev_out_csum = fp; gsp_phys_gr_prev_out_valid = true; } changed++; @@ -21546,8 +21869,7 @@ gr_page_unmapped: if (bi == GSP_PHYS_GR_DATA_BI_NONE) continue; wp = gsp_phys_chunk_page_va(dchunk, bi); - for (k = 0; k < GSP_PHYS_PAGE_SIZE; k += 4) - wc += get_unaligned_le32(wp + k); + (void)k; (void)wc; pr_info("vgsp-gr: window[%u] reloc=0x%llx first=%08x %08x %08x %08x changed=%u\n", bi, req->data[wi].reloc_va, @@ -21555,7 +21877,9 @@ gr_page_unmapped: get_unaligned_le32(wp + 4), get_unaligned_le32(wp + 8), get_unaligned_le32(wp + 12), - wc != bounce_csum_in[bi]); + gsp_phys_gr_fnv1a64( + wp, GSP_PHYS_PAGE_SIZE) != + bounce_csum_in[bi]); } } pr_info("vgsp-gr: bounce copy-out: %u pages written back, %u changed by engine, %u unchanged-skipped, %u alias-skipped, %u stale-skipped (first bi: %u %u %u %u %u %u %u %u)\n", @@ -21578,7 +21902,7 @@ gr_page_unmapped: for (gi = 0; gi < gsp_phys_gr_gap_n; gi++) { u32 bi = gsp_phys_gr_gap[gi].bi; const u8 *src = gsp_phys_chunk_page_va(dchunk, bi); - u32 ci, csum = 0; + u64 fp; if (!src) continue; @@ -21591,9 +21915,8 @@ gr_page_unmapped: gskipped++; continue; } - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(src + ci); - if (csum == bounce_csum_in[bi]) { + fp = gsp_phys_gr_fnv1a64(src, GSP_PHYS_PAGE_SIZE); + if (fp == bounce_csum_in[bi]) { /* Engine did not touch it: nothing to ship. */ gunchanged++; continue; @@ -21662,8 +21985,20 @@ gr_page_unmapped: jiffies_to_msecs(jiffies - start), gsp_phys_gr_ok); /* m128: post-success qmdwin words so a probe that stores a value into * the qmdwin page (e.g. p6's constant-read dump at w48) is readable - * from the success log (the post-mortem only prints on timeout). */ - if (boot_stage.rmcall_gr_qmdwin.pages) { + * from the success log (the post-mortem only prints on timeout). + * r434: read the launch's OWN bank (qmd_dst_mem) when available - + * the old bank-0-only read showed a stale page for odd epochs. */ + if (qmd_dst_mem) { + const u32 *qw = (const u32 *)qmd_dst_mem; + + pr_info("vgsp-gr: post-OK qmdwin w48..59=%08x %08x %08x %08x %08x %08x %08x %08x %08x %08x %08x %08x\n", + get_unaligned_le32(qw + 48), get_unaligned_le32(qw + 49), + get_unaligned_le32(qw + 50), get_unaligned_le32(qw + 51), + get_unaligned_le32(qw + 52), get_unaligned_le32(qw + 53), + get_unaligned_le32(qw + 54), get_unaligned_le32(qw + 55), + get_unaligned_le32(qw + 56), get_unaligned_le32(qw + 57), + get_unaligned_le32(qw + 58), get_unaligned_le32(qw + 59)); + } else if (boot_stage.rmcall_gr_qmdwin.pages) { const u32 *qw = (const u32 *)gsp_phys_chunk_page_va( &boot_stage.rmcall_gr_qmdwin, 0); diff --git a/kernel/gsp_phys_gpu.c b/kernel/gsp_phys_gpu.c index b93c254..74062dd 100644 --- a/kernel/gsp_phys_gpu.c +++ b/kernel/gsp_phys_gpu.c @@ -664,6 +664,29 @@ int gsp_phys_rmcall_gr_pcas_value(void) EXPORT_SYMBOL_GPL(gsp_phys_rmcall_gr_pcas_value); /* + * r434 CWD-slot-count path (gsp_phys_rmcall_gr_cwd_slot; default -1 = off). + * On Volta+ the SET_CWD_SLOT_COUNT method is gated by the falcon PRI bit + * NV_PGRAPH_PRI_SKED_DEBUG_1.ENABLE_CWD_SLOT_COUNT (the CUDA driver's + * turingLaunchSetCwdSlotCount: a masked PRI write via SET_MME_SHADOW_SCRATCH + * 0/1/2 + SET_FALCON04, then the method itself). When >= 0 the init preamble + * emits that PRI write (ENABLE_CWD_SLOT_COUNT=1) followed by + * SET_CWD_SLOT_COUNT with the given value, so the wedge's "grid truncated to + * ~4 CTAs" signature can be A/B'd against the CWD slot-count watermark. + * Read per launch; -1 restores the historical preamble byte-for-byte. + */ +static int gsp_phys_rmcall_gr_cwd_slot_param = -1; +module_param_named(gsp_phys_rmcall_gr_cwd_slot, + gsp_phys_rmcall_gr_cwd_slot_param, int, 0644); +MODULE_PARM_DESC(gsp_phys_rmcall_gr_cwd_slot, + "r434: enable the CWD slot count PRI (SKED_DEBUG_1 bit11) and emit SET_CWD_SLOT_COUNT with this value in the GR init preamble (-1 = off, default)"); + +int gsp_phys_rmcall_gr_cwd_slot_value(void) +{ + return gsp_phys_rmcall_gr_cwd_slot_param; +} +EXPORT_SYMBOL_GPL(gsp_phys_rmcall_gr_cwd_slot_value); + +/* * Two-step PCAS gate (default 0). With the explicit PCAS pair on, setting * this marks the QMD's REQUIRE_SCHEDULING_PCAS (w4 bit12) and emits the copy * PCAS (INVALIDATE only) followed by a schedule-only PCAS (SCHEDULE), the @@ -711,20 +734,30 @@ int gsp_phys_rmcall_gr_qmdwin_value(void) EXPORT_SYMBOL_GPL(gsp_phys_rmcall_gr_qmdwin_value); /* - * m64 forced module-semaphore release (gsp_phys_rmcall_gr_rel_mod; default 0). - * When set, a forwarded GR launch ignores the resolved guest release page and - * points QMD release0 at the module window-0 semaphore VA - * (GSP_PHYS_RMCALL_VA_SEM), which the host polls at the SEM chunk's host VA. - * Discriminant for "engine completed but the shadow release write was lost" - * vs "engine never completed": with the module release, a completed QMD is - * always observable. Also makes the existing module-release fallback - * observable (it previously set no poll source, so it could never complete). - */ -static int gsp_phys_rmcall_gr_rel_mod_param; + * m64 forced module-semaphore release (gsp_phys_rmcall_gr_rel_mod). + * + * DEFAULT 1 since r453. When set, a forwarded GR launch ignores the resolved + * guest release page and points QMD release0 at the module window-0 semaphore + * VA (GSP_PHYS_RMCALL_VA_SEM), which the host polls at the SEM chunk's host + * VA; the GUEST-VISIBLE release word is authored by the host only on the + * completion edge - after the retire wait and after the bounce copy-out has + * shipped the engine's stores to the guest pages. Evidence (r453, the FNV + * change-detection fix on top): + * - with 0: the guest's release could be satisfied by the engine's own + * early write while the host copy-out was still pending, so a client's + * stream-ordered readback (streams2's c2, memcpy_roundtrip's memset + * phase, streams_events' chain) read PRE-launch content: 11/15 battery. + * - with 1: 13/15 battery (streams_events and memcpy_roundtrip flipped + * green), no stalls across the full battery, mail+saxpy clean. + * The r414 stall hole (a launch that never reached the hold edge left the + * guest word unauthored) is closed by the r420b per-outcome authoring path; + * one-line rollback: echo 0 > .../parameters/gsp_phys_rmcall_gr_rel_mod. + */ +static int gsp_phys_rmcall_gr_rel_mod_param = 1; module_param_named(gsp_phys_rmcall_gr_rel_mod, gsp_phys_rmcall_gr_rel_mod_param, int, 0644); MODULE_PARM_DESC(gsp_phys_rmcall_gr_rel_mod, - "Force the forwarded GR QMD release0 to the module semaphore (default 0; 1 = m64 engine-completion discriminant)"); + "Force the forwarded GR QMD release0 to the module semaphore; guest release authored only on the post-copy-out completion edge (default 1 since r453; 0 = historical engine-authored guest release)"); int gsp_phys_rmcall_gr_rel_mod_value(void) { @@ -1057,6 +1090,33 @@ int gsp_phys_gr_icache_inv_value(void) EXPORT_SYMBOL_GPL(gsp_phys_gr_icache_inv_value); /* + * r433 RELEASE_MEMBAR_TYPE A/B (gsp_phys_gr_rel_membar; default 0). The + * captured guest QMDs carry RELEASE_MEMBAR_TYPE=FE_NONE (QMDV00_06 word 11 + * bit 14 clear, w11=0x44010000), so when the front-end completes a grid it + * writes the QMD release semaphore without first flushing its pending + * system-memory stores. The post-dispatch store-loss signature (the engine + * processes the QMD, writes the release, the grid completes and ZERO data + * pages are host-visible) is exactly what an unflushed write-back L2 looks + * like from the host: the release is host-visible while the data sits in the + * GPU's L2. Forcing RELEASE_MEMBAR_TYPE=FE_SYSMEMBAR makes the release + * ordering guarantee the stores are in memory before the semaphore write, so + * the host's retire scan (and the copy-out) can never observe a completion + * whose data has not landed. 0 keeps the guest's captured value (the + * historical behaviour) for a one-flip A/B; read per launch. + */ +static int gsp_phys_gr_rel_membar_param; +module_param_named(gsp_phys_gr_rel_membar, gsp_phys_gr_rel_membar_param, + int, 0644); +MODULE_PARM_DESC(gsp_phys_gr_rel_membar, + "r433: force QMD RELEASE_MEMBAR_TYPE=FE_SYSMEMBAR (word11 bit14) on forwarded GR launches (default 0 = guest's captured FE_NONE)"); + +int gsp_phys_gr_rel_membar_value(void) +{ + return gsp_phys_gr_rel_membar_param; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_rel_membar_value); + +/* * m179 no-write retry budget (gsp_phys_gr_nowrite_retries; default 3, clamped * to [0, 4]). The m153 outcome retry loop re-executes a data-bearing launch * whose retirement wait saw zero engine writes; it shared the timeout ladder's @@ -1124,6 +1184,60 @@ int gsp_phys_gr_nowrite_retries_value(void) EXPORT_SYMBOL_GPL(gsp_phys_gr_nowrite_retries_value); /* + * r445 wedge hard re-arm (gsp_phys_gr_wedge_rearm; default 0 = off). When the + * outcome-retry's warm GR semaphore probe proves the engine wedged for the + * context (the proven-sticky post-dispatch no-store state), this knob makes + * the submit attempt ONE bounded recovery instead of breaking immediately: + * a TSG PREEMPT(bWait=1) forces the FECS to save/unload the channel's + * context, then the mode-35 lane-enable tail + GPFIFO schedules are re-issued + * (the m155 extended re-arm) and the semaphore probe is repeated. A pass + * means the lane revived for this launch and the normal retry ladder + * continues; a fail leaves the historical WEDGED break. Read once per wedge + * episode. Off by default for A/B. + */ +static int gsp_phys_gr_wedge_rearm_param; +module_param_named(gsp_phys_gr_wedge_rearm, gsp_phys_gr_wedge_rearm_param, + int, 0644); +MODULE_PARM_DESC(gsp_phys_gr_wedge_rearm, + "r445: on the proven engine-wedged sem probe, attempt a bounded preempt+resched hard re-arm before giving up (default 0=off)"); + +int gsp_phys_gr_wedge_rearm_value(void) +{ + return gsp_phys_gr_wedge_rearm_param ? 1 : 0; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_wedge_rearm_value); + +/* + * r449 SHADER-CONSTANT-CACHE invalidate (gsp_phys_gr_ccache_inv; default 0). + * The forwarded QMD's cache-control block (QMDV02_03 word 5) carries the + * per-cache invalidate bits; the captured guest QMDs always set bits 26..29 + * (texture header/sampler/data + shader DATA) and the M91 fix adds bit 30 + * (INSTRUCTION), but bit 31 (SHADER CONSTANT - where a kernel's launch + * parameters live) is never set. The forwarded launch's param page is + * rewritten in place between launches (the two-bank win0 pool), so a stale + * constant-cache line can serve the PREVIOUS launch's parameters - the + * kernel then computes with the previous launch's pointers/sizes and the + * current launch's tracked output pages stay untouched ("no stores"). + * Mode: 0 = off, 1 = set bit 31 only, 2 = force the full 0xff000000 + * invalidate set. Read at submission time, so it can be A/B'd at runtime. + */ +static int gsp_phys_gr_ccache_inv_param; +module_param_named(gsp_phys_gr_ccache_inv, gsp_phys_gr_ccache_inv_param, + int, 0644); +MODULE_PARM_DESC(gsp_phys_gr_ccache_inv, + "r449: forwarded-QMD shader-constant-cache invalidate: 0=off, 1=bit31, 2=all invalidate bits (default 0)"); + +unsigned int gsp_phys_gr_ccache_inv_value(void) +{ + if (gsp_phys_gr_ccache_inv_param < 0) + return 0; + if (gsp_phys_gr_ccache_inv_param > 2) + return 2; + return (unsigned int)gsp_phys_gr_ccache_inv_param; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_ccache_inv_value); + +/* * r410 wedge quiet window (gsp_phys_gr_wedge_quiet_ms; default 40000, clamped * [0,120000]; 0 disables). The sticky per-lane no-execute wedge (100% of * launches no-write while kept busy) is proven to CLEAR when the lane stays diff --git a/kernel/gsp_phys_gpu.h b/kernel/gsp_phys_gpu.h index e02d6d8..36c2220 100644 --- a/kernel/gsp_phys_gpu.h +++ b/kernel/gsp_phys_gpu.h @@ -388,10 +388,22 @@ unsigned int gsp_phys_gr_first_poll_ms_value(void); * (default 1; 0 = pre-m91 QMD words). */ int gsp_phys_gr_icache_inv_value(void); +/* r433: force QMD RELEASE_MEMBAR_TYPE=FE_SYSMEMBAR (QMDV00_06 word11 bit14) + * on forwarded GR launches so the engine flushes pending sysmem stores + * before writing the release semaphore (default 0 = guest's FE_NONE). */ +int gsp_phys_gr_rel_membar_value(void); + +/* r434: enable the CWD slot count PRI (SKED_DEBUG_1 bit11) and emit + * SET_CWD_SLOT_COUNT with this value in the GR init preamble + * (-1 = off, default). */ +int gsp_phys_rmcall_gr_cwd_slot_value(void); + /* m179: dedicated re-execution budget for a data-bearing GR launch whose * retirement wait saw zero engine writes, independent of the timeout ladder's * attempt counter (default 3, clamped [0,4]). */ int gsp_phys_gr_nowrite_retries_value(void); +int gsp_phys_gr_wedge_rearm_value(void); +unsigned int gsp_phys_gr_ccache_inv_value(void); /* STALLFREE: absolute per-launch GR submit wall-clock budget in ms, clamped * [10000,600000]; 0 = only the built-in per-ladder bounds (default 60000). diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 931669c..28e5906 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -486,6 +486,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 +825,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 +1606,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 +1967,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; @@ -4820,7 +4875,8 @@ static int svc_fill_1d(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u64 *done_out); static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, - u32 src_phys, u32 src_aperture, u64 *done_out); + u32 src_phys, u32 src_aperture, u64 *done_out, + u32 owner_ch, u32 owner_idx, u64 owner_hash); /* * r402 row-wise wrapper: a 2D (pitched) fill runs the 1D fill once per @@ -5938,12 +5994,29 @@ static void svc_copy_audit_chunk(struct nv_vgpu_device *vgpu, u64 dst_va, struct svc_copy_snap { u64 dst, src; size_t total; + /* + * Owner identity (r444): the channel, ring entry and push-content + * hash of the decode pass that captured this snapshot. A snapshot + * exists to serve ONE entry's retried copy (the held staging read); + * keying it by (dst,src,total) alone let a snapshot captured for an + * earlier entry poison a later copy with the same operand triple - + * the wedge_ab "bad = mutated-element count" class: step N's H2D + * executed from step N-1's snapshot because the re-decode of the + * completed step N-1 entry re-captured its (unchanged) source after + * its copy had already run, and step N's copy then consumed it. + * The pass hash changes whenever the guest rewrites the push, so a + * ring slot reused by a fresh push can never match a stale owner. + */ + u64 owner_hash; + u32 owner_ch; + u32 owner_idx; u8 *buf; unsigned long stamp; }; static struct svc_copy_snap svc_copy_snaps[SVC_COPY_SNAP_N]; -static u8 *svc_copy_snap_lookup(u64 dst, u64 src, size_t total) +static u8 *svc_copy_snap_lookup(u64 dst, u64 src, size_t total, + u32 owner_ch, u32 owner_idx, u64 owner_hash) { u32 i; @@ -5951,12 +6024,16 @@ static u8 *svc_copy_snap_lookup(u64 dst, u64 src, size_t total) if (svc_copy_snaps[i].buf && svc_copy_snaps[i].dst == dst && svc_copy_snaps[i].src == src && - svc_copy_snaps[i].total == total) + svc_copy_snaps[i].total == total && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx && + svc_copy_snaps[i].owner_hash == owner_hash) return svc_copy_snaps[i].buf; return NULL; } -static void svc_copy_snap_free(u64 dst, u64 src, size_t total) +static void svc_copy_snap_free(u64 dst, u64 src, size_t total, + u32 owner_ch, u32 owner_idx, u64 owner_hash) { u32 i; @@ -5964,16 +6041,45 @@ static void svc_copy_snap_free(u64 dst, u64 src, size_t total) if (svc_copy_snaps[i].buf && svc_copy_snaps[i].dst == dst && svc_copy_snaps[i].src == src && - svc_copy_snaps[i].total == total) { + svc_copy_snaps[i].total == total && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx && + svc_copy_snaps[i].owner_hash == owner_hash) { kfree(svc_copy_snaps[i].buf); svc_copy_snaps[i].buf = NULL; return; } } +/* + * 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 __maybe_unused svc_copy_snap_free_owner(u32 owner_ch, u32 owner_idx) +{ + u32 i; + + for (i = 0; i < SVC_COPY_SNAP_N; i++) + if (svc_copy_snaps[i].buf && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx) { + kfree(svc_copy_snaps[i].buf); + svc_copy_snaps[i].buf = NULL; + } +} + /* Capture the source of a virtual-operand copy about to be held. */ static void svc_copy_snap_capture(struct nv_vgpu_device *vgpu, u64 dst, - u64 src, size_t total) + u64 src, size_t total, u32 owner_ch, + u32 owner_idx, u64 owner_hash) { struct svc_copy_snap *s = NULL; u8 *buf; @@ -5982,8 +6088,27 @@ static void svc_copy_snap_capture(struct nv_vgpu_device *vgpu, u64 dst, if (!total || total > SVC_COPY_SNAP_MAX) return; - if (svc_copy_snap_lookup(dst, src, total)) + if (svc_copy_snap_lookup(dst, src, total, owner_ch, owner_idx, + owner_hash)) return; /* already captured */ + /* + * r456: evict any snapshot left by the SAME ring slot under a + * DIFFERENT push hash (the slot was reused by a fresh push). The + * surviving-across-retries snapshot (svc_opdef_clear APPLY no + * longer frees it) must never be readable by the next occupant of + * the slot; the pass hash changes whenever the guest rewrites the + * push, so a same-slot different-hash snapshot is by definition + * stale for the new push. + */ + for (i = 0; i < SVC_COPY_SNAP_N; i++) { + if (svc_copy_snaps[i].buf && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx && + svc_copy_snaps[i].owner_hash != owner_hash) { + kfree(svc_copy_snaps[i].buf); + svc_copy_snaps[i].buf = NULL; + } + } for (i = 0; i < SVC_COPY_SNAP_N; i++) { if (!svc_copy_snaps[i].buf) { s = &svc_copy_snaps[i]; @@ -6019,10 +6144,14 @@ static void svc_copy_snap_capture(struct nv_vgpu_device *vgpu, u64 dst, s->dst = dst; s->src = src; s->total = total; + s->owner_ch = owner_ch; + s->owner_idx = owner_idx; + s->owner_hash = owner_hash; s->stamp = jiffies; s->buf = buf; - pr_info_ratelimited("vgsp-sema: copy snapshot captured dst=0x%llx src=0x%llx total=%zu first4=%08x %08x %08x %08x\n", - dst, src, total, + pr_info_ratelimited("vgsp-sema: copy snapshot captured dst=0x%llx src=0x%llx total=%zu owner=0x%x/%u/%016llx first4=%08x %08x %08x %08x\n", + dst, src, total, owner_ch, owner_idx, + owner_hash, total >= 4 ? *(const u32 *)buf : 0, total >= 8 ? *(const u32 *)(buf + 4) : 0, total >= 12 ? *(const u32 *)(buf + 8) : 0, @@ -6039,7 +6168,8 @@ static const char *svc_gr_src_probe(struct nv_vgpu_device *vgpu, u64 va, u64 *pa_out, u32 *ap_out); static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, u32 src_phys, u32 src_aperture, - u64 *done_out, u32 rows, u32 dpitch, u32 spitch) + u64 *done_out, u32 rows, u32 dpitch, u32 spitch, + u32 owner_ch, u32 owner_idx, u64 owner_hash) { if (rows > 1 && dpitch) { u64 done = 0; @@ -6052,7 +6182,8 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, rc = svc_copy_1d(vgpu, dst + (u64)r * dpitch, src + (u64)r * spitch, total, dst_phys, dst_aperture, src_phys, - src_aperture, &row_done); + src_aperture, &row_done, + owner_ch, owner_idx, owner_hash); done += row_done; if (rc) break; @@ -6062,12 +6193,13 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, return rc; } return svc_copy_1d(vgpu, dst, src, total, dst_phys, dst_aperture, - src_phys, src_aperture, done_out); + src_phys, src_aperture, done_out, + owner_ch, owner_idx, owner_hash); } static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, u32 src_phys, u32 src_aperture, - u64 *done_out) + u64 *done_out, u32 owner_ch, u32 owner_idx, u64 owner_hash) { u8 bounce[256]; size_t o; @@ -6098,7 +6230,8 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot svc.cur_ch, dst, src, total, dst_phys, dst_aperture, src_phys, src_aperture); } - snap = svc_copy_snap_lookup(dst_orig, src_orig, total); + snap = svc_copy_snap_lookup(dst_orig, src_orig, total, owner_ch, + owner_idx, owner_hash); if (!dst_phys) { u64 pa; @@ -6232,7 +6365,8 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot */ if (!src_phys && src_orig >= SVC_VAS_FLOOR) svc_copy_snap_capture(vgpu, dst_orig, src_orig, - total); + total, owner_ch, owner_idx, + owner_hash); return -ENOENT; } /* @@ -6261,7 +6395,9 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot dst_orig, dst_prov, ipa); if (!src_phys && src_orig >= SVC_VAS_FLOOR) svc_copy_snap_capture(vgpu, dst_orig, - src_orig, total); + src_orig, total, + owner_ch, owner_idx, + owner_hash); return -ENOENT; } } @@ -7036,7 +7172,8 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot svc_ledger_store(dst_orig + po, dst + po, cap); } if (snap) - svc_copy_snap_free(dst_orig, src_orig, total); + svc_copy_snap_free(dst_orig, src_orig, total, owner_ch, + owner_idx, owner_hash); if (done_out) *done_out = (u64)o; /* o == total on the success path */ return 0; @@ -7238,7 +7375,8 @@ static struct svc_wait_hold *svc_wait_hold_get(u32 h, bool create) } static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops, - u32 n_ops, u32 ch_h, u64 *ops_done, u32 ops_base) + u32 n_ops, u32 ch_h, u64 *ops_done, u32 ops_base, + u32 entry_idx, u64 pass_hash) { u32 i; static u64 bigop_log; @@ -7391,6 +7529,18 @@ svc_wait_done: */ if (ops_done && seq < 64 && ((*ops_done >> seq) & 1)) { svc.ops_once_skipped++; + /* + * r456: this op already ran in an earlier pass. Any + * staging snapshot a later re-decode pass re-captured + * for it (the capture sites run on every decode) is + * a leak that the r444 APPLY-free used to sweep; with + * that free removed (r456 keeps the snapshot alive + * across retries) drop it here so it can never serve + * a re-push of the ring slot. + */ + if (m->type == 1 && m->bytes <= SVC_COPY_SNAP_MAX) + svc_copy_snap_free(m->dst, m->src, m->bytes, + ch_h, entry_idx, pass_hash); continue; } @@ -7404,7 +7554,8 @@ svc_wait_done: m->phys, m->aperture, m->src_phys, m->src_aperture, &ce_done, m->rows, m->dpitch, - m->spitch); + m->spitch, ch_h, entry_idx, + pass_hash); else rc = svc_fill(vgpu, m->dst, m->value, m->elem, m->count, m->phys, m->aperture, @@ -9385,7 +9536,8 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, * REPORT_SEM comment below. */ if (allow_ops && n_ops && !gsp_sema_skip_ops) { svc_run_ops(vgpu, ops, n_ops, ch->h_object, - ops_done, ops_seq - n_ops); + ops_done, ops_seq - n_ops, + entry_idx, svc_pass_hash); svc.ring_ops++; ops_runs++; } @@ -9585,7 +9737,8 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, */ if (allow_ops && n_ops && !gsp_sema_skip_ops) { svc_run_ops(vgpu, ops, n_ops, ch->h_object, - ops_done, ops_seq - n_ops); + ops_done, ops_seq - n_ops, + entry_idx, svc_pass_hash); svc.ring_ops++; ops_runs++; } @@ -9809,6 +9962,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: @@ -10116,6 +10286,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) { @@ -10295,7 +10498,10 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, svc_copy_snap_capture(vgpu, st.offset_out, st.offset_in, - st.line_length); + st.line_length, + ch->h_object, + entry_idx, + svc_pass_hash); } /* * r308: the `!replay` gate above kept this @@ -10386,7 +10592,10 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, svc_copy_snap_capture(vgpu, lm->dst, lm->src, - lm->bytes); + lm->bytes, + ch->h_object, + entry_idx, + svc_pass_hash); if (gsp_sema_dbg && svc.snap_gate_log < 256) { svc.snap_gate_log++; @@ -10707,7 +10916,8 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, gsp_sema_skip_ops); } svc_run_ops(vgpu, ops, n_ops, ch->h_object, - ops_done, ops_seq - n_ops); + ops_done, ops_seq - n_ops, + entry_idx, svc_pass_hash); svc.ring_ops++; ops_runs++; } @@ -15805,6 +16015,29 @@ static void svc_opdef_clear(struct svc_opdef_ent *e, u32 caller) e->n ? e->rel[0].va : 0ull, e->n ? e->rel[0].payload : 0u); } + /* + * 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. + */ memset(e, 0, sizeof(*e)); } @@ -21234,10 +21467,55 @@ 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; + } + } 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 == @@ -21289,6 +21567,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 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/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/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/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/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/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/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/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/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 c150692..4f1430f 100755 --- a/tools/cuda-ab/soak.sh +++ b/tools/cuda-ab/soak.sh @@ -192,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/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_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/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]) 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 |