summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--docs/03-architecture/13-vexecution-plane.md33
-rw-r--r--kernel/gsp_phys_booter.c205
-rw-r--r--kernel/gsp_phys_gpu.c62
-rw-r--r--kernel/gsp_phys_gpu.h1
-rw-r--r--tools/cuda-ab/ip_probe.cu128
-rw-r--r--tools/cuda-ab/ip_probe2.cu201
-rw-r--r--tools/cuda-ab/ip_probe3.cu142
-rw-r--r--tools/cuda-ab/rt_probe.cu165
-rwxr-xr-xtools/cuda-ab/run/ip4bin0 -> 825200 bytes
-rwxr-xr-xtools/cuda-ab/run/ip5bin0 -> 829288 bytes
-rwxr-xr-xtools/cuda-ab/run/ip6bin0 -> 829376 bytes
-rwxr-xr-xtools/cuda-ab/run/ip_probebin0 -> 829288 bytes
-rwxr-xr-xtools/cuda-ab/run/ip_probe2bin0 -> 833680 bytes
-rwxr-xr-xtools/cuda-ab/run/ip_probe3bin0 -> 829296 bytes
-rwxr-xr-xtools/cuda-ab/run/rt_probebin0 -> 825184 bytes
-rwxr-xr-xtools/cuda-ab/run/scale_probebin0 -> 833704 bytes
-rwxr-xr-xtools/cuda-ab/run/streams2_probebin0 -> 833592 bytes
-rw-r--r--tools/cuda-ab/scale_probe.cu245
-rw-r--r--tools/cuda-ab/streams2_probe.cu299
19 files changed, 1433 insertions, 48 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/kernel/gsp_phys_booter.c b/kernel/gsp_phys_booter.c
index 967dc25..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
@@ -16472,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
@@ -17389,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
@@ -17838,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();
@@ -17884,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)
@@ -17928,7 +17986,7 @@ static void gsp_phys_gr_qmd64_post(const u8 *deposit, const char *tag)
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;
@@ -18082,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;
@@ -18160,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
@@ -18205,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 {
@@ -18242,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;
@@ -18291,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;
@@ -18299,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;
@@ -18309,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();
@@ -18574,6 +18649,7 @@ 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,
@@ -18906,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
@@ -19084,6 +19160,32 @@ 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
@@ -21414,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);
@@ -21518,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)
@@ -21539,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
@@ -21550,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;
}
{
@@ -21600,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());
@@ -21621,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;
@@ -21657,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++;
@@ -21738,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,
@@ -21747,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",
@@ -21770,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;
@@ -21783,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;
diff --git a/kernel/gsp_phys_gpu.c b/kernel/gsp_phys_gpu.c
index b61929e..74062dd 100644
--- a/kernel/gsp_phys_gpu.c
+++ b/kernel/gsp_phys_gpu.c
@@ -734,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)
{
@@ -1198,6 +1208,36 @@ int gsp_phys_gr_wedge_rearm_value(void)
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 46d841d..36c2220 100644
--- a/kernel/gsp_phys_gpu.h
+++ b/kernel/gsp_phys_gpu.h
@@ -403,6 +403,7 @@ int gsp_phys_rmcall_gr_cwd_slot_value(void);
* 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/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/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
new file mode 100755
index 0000000..54937e6
--- /dev/null
+++ b/tools/cuda-ab/run/ip4
Binary files differ
diff --git a/tools/cuda-ab/run/ip5 b/tools/cuda-ab/run/ip5
new file mode 100755
index 0000000..c4b4a39
--- /dev/null
+++ b/tools/cuda-ab/run/ip5
Binary files differ
diff --git a/tools/cuda-ab/run/ip6 b/tools/cuda-ab/run/ip6
new file mode 100755
index 0000000..56815c2
--- /dev/null
+++ b/tools/cuda-ab/run/ip6
Binary files differ
diff --git a/tools/cuda-ab/run/ip_probe b/tools/cuda-ab/run/ip_probe
new file mode 100755
index 0000000..08fb86e
--- /dev/null
+++ b/tools/cuda-ab/run/ip_probe
Binary files differ
diff --git a/tools/cuda-ab/run/ip_probe2 b/tools/cuda-ab/run/ip_probe2
new file mode 100755
index 0000000..dfe6ab9
--- /dev/null
+++ b/tools/cuda-ab/run/ip_probe2
Binary files differ
diff --git a/tools/cuda-ab/run/ip_probe3 b/tools/cuda-ab/run/ip_probe3
new file mode 100755
index 0000000..8a34c05
--- /dev/null
+++ b/tools/cuda-ab/run/ip_probe3
Binary files differ
diff --git a/tools/cuda-ab/run/rt_probe b/tools/cuda-ab/run/rt_probe
new file mode 100755
index 0000000..48d866f
--- /dev/null
+++ b/tools/cuda-ab/run/rt_probe
Binary files differ
diff --git a/tools/cuda-ab/run/scale_probe b/tools/cuda-ab/run/scale_probe
new file mode 100755
index 0000000..ca3fc20
--- /dev/null
+++ b/tools/cuda-ab/run/scale_probe
Binary files differ
diff --git a/tools/cuda-ab/run/streams2_probe b/tools/cuda-ab/run/streams2_probe
new file mode 100755
index 0000000..206d6d6
--- /dev/null
+++ b/tools/cuda-ab/run/streams2_probe
Binary files differ
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/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;
+}
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-22 09:22:16 +0000