vgsp: fail4 r226 - zero-filter alias witnesses + GRID-WIT QMD geometry dump

r226a: the green victory_mail run produced 10 IN-ALIAS-WIT lines that are
all false positives - consecutive all-zero pages from fresh cudaMalloc
runs (csum=0x00000000), not aliasing evidence. Filter all-zero pages out of
both alias comparisons: a page counts as zero only when the csum is zero AND
every word is zero (csum alone can hide cancelling contents); zero pages are
neither candidates nor partners.

Also: the r224 band corruption (rows >= 16 of every launch replicate one
tile) starts at a CTA-band boundary while the QMD is passed through
verbatim - but the current dumps never show the geometry words the engine
actually fetches. GRID-WIT logs, once per launch, the geometry-relevant QMD
words of the deposit the engine will run: w12/w13 = CTA raster
width/height/depth (gridDim), w18/w19 = CTA thread dimensions (blockDim),
w18.lo = QMD version nibbles, plus w4/w5/w8/w17 (PCAS/cache-control,
prefetch low word, shared-memory size). With the guest-side GEO line added
to transformer_mt.cu (8 lines per run, one per chain) the next capture can
prove or refute an emission mismatch in one comparison.

And IN-ALIAS-WIT now prints both HPAs with a classification token: equal
hpa = the resolver served one page for several VAs; distinct hpa with equal
contents = the guest pages genuinely hold the same bytes (upstream write
aliasing). Logging-only; no behaviour change, safe for the green clients.

(cherry picked from commit dc1e693d598d3a32d1e6a87b94ff91dcbc99101c)

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;