diff options
Diffstat (limited to 'kernel/gsp_phys_booter.c')
| -rw-r--r-- | kernel/gsp_phys_booter.c | 79 |
1 files changed, 69 insertions, 10 deletions
diff --git a/kernel/gsp_phys_booter.c b/kernel/gsp_phys_booter.c index bc0ca74..e336247 100644 --- a/kernel/gsp_phys_booter.c +++ b/kernel/gsp_phys_booter.c @@ -17737,6 +17737,26 @@ static void gsp_phys_gr_page_spot_check(const u8 *dst, const u8 *src, } /* + * 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 + * the IN-ALIAS-WIT false positives of the r226a green run (csum=0x00000000 + * sequences). A page counts as zero only when its csum is zero AND every + * word is zero (csum==0 alone can hide cancelling contents). + */ +static inline bool gsp_phys_gr_page_all_zero(u32 csum, const u8 *p) +{ + u32 i; + + if (csum) + return false; + for (i = 0; i < GSP_PHYS_PAGE_SIZE; i += 4) + if (get_unaligned_le32(p + i)) + return false; + return true; +} + +/* * m153: fill the bounce window from the request's guest pages and capture a * per-page checksum baseline. Used before the first publish and again before * every outcome retry (a mis-executing first attempt can leave stale or @@ -17748,7 +17768,7 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, { struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); u32 di, copied = 0; - u64 prev_va = 0; + u64 prev_va = 0, prev_hpa = 0; u32 prev_csum = 0; for (di = 0; di < req->n_data; di++) { @@ -17756,6 +17776,7 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, const u8 *src; u8 *dst; u32 ci, csum = 0; + bool zero; if (bi == GSP_PHYS_GR_DATA_BI_NONE || !req->data[di].hpa) continue; @@ -17769,20 +17790,35 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) csum += get_unaligned_le32(dst + ci); csum_in[bi] = csum; + zero = gsp_phys_gr_page_all_zero(csum, dst); /* * IN-ALIAS-WIT (COPYWIT r224 follow-up): two consecutive tracked * recs whose guest content is identical prove a replicated INPUT - * region (one page served for many VAs). Logged with both VAs so - * the source can be named; bounded boot-wide. + * region (one page served for many VAs). The HPAs tell the two + * classes apart: equal hpa = the resolver served ONE page for + * several VAs (resolution aliasing); different hpa = the guest + * pages genuinely hold the same bytes (upstream write aliasing, + * e.g. an H2D that landed on one page). Bounded boot-wide; all- + * zero pages are excluded on both sides (r226: a fresh client's + * zero runs are not aliasing evidence). */ - if (prev_va && csum == prev_csum && gsp_phys_gr_copywit_enabled() && + if (prev_va && !zero && csum == prev_csum && + gsp_phys_gr_copywit_enabled() && gsp_phys_gr_in_alias_log < 12) { gsp_phys_gr_in_alias_log++; - pr_info_ratelimited("vgsp-gr: IN-ALIAS-WIT va=0x%llx prev_va=0x%llx csum=0x%08x (identical guest page content)\n", - req->data[di].va, prev_va, csum); + pr_info_ratelimited("vgsp-gr: IN-ALIAS-WIT va=0x%llx hpa=0x%llx prev_va=0x%llx prev_hpa=0x%llx csum=0x%08x (%s)\n", + req->data[di].va, + req->data[di].hpa, prev_va, + prev_hpa, csum, + prev_hpa == req->data[di].hpa ? + "resolver served one page for many VAs" : + "distinct pages, identical contents"); + } + if (!zero) { + prev_va = req->data[di].va; + prev_hpa = req->data[di].hpa; + prev_csum = csum; } - prev_va = req->data[di].va; - prev_csum = csum; copied++; } dma_wmb(); @@ -18629,6 +18665,25 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, } /* + * GRID-WIT (r226): the engine reads CTA geometry from this verbatim + * QMD copy, and the copy-out witnesses show the corruption of the + * failing clients starts exactly at a CTA-band boundary (rows >= 16 + * replicate one tile). Dump the geometry-relevant words once per + * launch so the next capture proves what raster/thread dimensions the + * engine actually sees and lets them be compared against the client's + * own launch config: w12 = CTA_RASTER_WIDTH (gridDim.x), + * w13.lo = CTA_RASTER_HEIGHT (gridDim.y), w13.hi = DEPTH (z), + * w14/w15 = raster resume W/H, w18.hi = CTA_THREAD_DIMENSION0 + * (blockDim.x), w19.lo/hi = DIM1/DIM2 (blockDim.y/z), w18.lo = + * QMD_VERSION/MAJOR nibbles. w4/w5/w8/w17: PCAS/cache-control, + * prefetch pair low word, shared-memory size. + */ + if (gsp_phys_gr_copywit_enabled()) + pr_info("vgsp-gr: GRID-WIT w4=%08x w5=%08x w8=%08x w12=%08x w13=%08x w14=%08x w15=%08x w17=%08x w18=%08x w19=%08x\n", + qmd[4], qmd[5], qmd[8], qmd[12], qmd[13], qmd[14], + qmd[15], qmd[17], qmd[18], qmd[19]); + + /* * Pages: Tier 1 maps the resolved HPA at the guest VA; a page without * a usable Tier-1 mapping falls to Tier 2 when captured inline bytes * cover it and a QMD pointer pair names it (the pair is rewritten to @@ -20685,9 +20740,13 @@ gr_page_unmapped: * OUT-ALIAS-WIT: identical first 8 bytes as * an earlier changed page of this launch = * the replicated-tile structure, with both - * recs named. + * recs named. r226: all-zero pages are + * excluded on both sides (the fresh-client + * fill stages zero runs; they are not alias + * evidence). */ - if (gsp_phys_gr_copywit_enabled()) { + if (gsp_phys_gr_copywit_enabled() && + !gsp_phys_gr_page_all_zero(csum, src)) { u32 wi; for (wi = 0; wi < gsp_phys_gr_outwit_n; |