summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--docs/03-architecture/06-host-module-and-hybrid-model.md33
-rw-r--r--docs/04-implementation/01-known-knowns-and-unknowns.md12
-rw-r--r--kernel/bar_emulation.c18
-rw-r--r--kernel/bar_init.c41
-rw-r--r--kernel/gsp_cuda_bridge.c495
-rw-r--r--kernel/gsp_cuda_bridge.h34
-rw-r--r--kernel/gsp_emulation.c13
-rw-r--r--kernel/gsp_rm_emulation.c18
-rw-r--r--kernel/gsp_rpc_client.c2
-rw-r--r--kernel/gsp_sema_service.c1579
-rw-r--r--kernel/gsp_sema_service.h7
-rw-r--r--kernel/include/nvidia-vgpu-vfio.h20
-rw-r--r--kernel/main.c31
-rwxr-xr-xtools/cuda-ab/run/vector_addbin820856 -> 820856 bytes
-rwxr-xr-xtools/cuda-ab/run/victory_mailbin829280 -> 829552 bytes
-rw-r--r--tools/cuda-ab/victory_mail.cu80
-rwxr-xr-xtools/cuda-bridge/cuda_bridgebin47840 -> 47936 bytes
-rw-r--r--tools/cuda-bridge/cuda_bridge.c92
18 files changed, 2182 insertions, 293 deletions
diff --git a/docs/03-architecture/06-host-module-and-hybrid-model.md b/docs/03-architecture/06-host-module-and-hybrid-model.md
index 7f670d7..dbf0c00 100644
--- a/docs/03-architecture/06-host-module-and-hybrid-model.md
+++ b/docs/03-architecture/06-host-module-and-hybrid-model.md
@@ -224,6 +224,38 @@ and the kernel-side bookkeeping records physical pages without a GPU VA.
Until a durable join is implemented, a session can feed verified pairs with
the `sema_map_add` debug parameter.
+### 4.4 Copy execution rules (correctness invariants)
+
+Two rules are load-bearing for any client-visible data transfer; both were
+established by live differential runs of a device-to-host/host-to-device
+pipeline:
+
+1. **A decoded copy runs at most once per GPFIFO entry.** The service pump
+ intentionally replays recently consumed ring entries when a new batch of
+ translations is injected (so releases that previously missed a
+ translation can land). Release application is forward-progress guarded,
+ but *copy execution is not idempotent*: a CE copy may read a client
+ staging buffer that the client has since refilled for the next chunk, so
+ re-running a consumed copy overwrites finished data with the next chunk's
+ bytes (observed as a deterministic corruption pattern in the transferred
+ buffer). The pump therefore records the raw GPFIFO entry words
+ (`e0 | e1 << 32`) per slot and, when a replayed entry matches the recorded
+ identity, suppresses its operations while still processing releases. A
+ slot whose entry words change (a recycled slot in a later epoch) executes
+ normally.
+2. **A virtual operand is translated per 4 KB page.** A GPU VA range is a
+ list of independent page mappings; resolving only the first page and
+ walking the physical address linearly is correct only while the mapping
+ is physically contiguous. Client staging buffers are frequently split
+ (a 1 MB range was observed as 195 contiguous pages plus 61 pages
+ elsewhere), so a linear walk reads unrelated memory past the split. Every
+ multi-page copy translates each page through the same trusted resolver
+ (pins, learned map, published page directories, targeted resolve) before
+ moving its 256-byte bounce.
+
+The same page-granular discipline already applies to the host-side bridge
+read/write helpers and to semaphore I/O.
+
## 5. The forwarding path (design)
```
@@ -332,6 +364,7 @@ in `kernel-open` first).
| Message queues + RPC + all init controls observed from stock 570.144 | implemented |
| Guest sysmem access + VA↔PA discovery | implemented (debug-grade, to be productized) |
| Push decoding + semaphore completion | demonstrated live; being productized |
+| Emulation-mode CUDA end-to-end (host bridge executes client kernels on the physical GPU) | verified live and repeatable: mail kernel bytes authored by GPU execution and read back, `vectorAdd` over 1 M elements returning `bad=0` and the reference checksum, deterministic across consecutive supervised cycles (see `docs/04-implementation/01-known-knowns-and-unknowns.md`) |
| Guest-RAM physical path | implemented default-off: guest RAM is file-backed shared memory and GPA→HPA goes through its page cache, verified live on the CE |
| Page-table population for server-owned ranges | implemented default-off: forwarded CE copies map resolved host pages through module-built GMMU scratch page tables; real guest page tables remain |
| GPU work forwarding (GR/CE execution) | CE forwarding, engine-authored release relay and fault relay implemented default-off and verified live; GR execution open |
diff --git a/docs/04-implementation/01-known-knowns-and-unknowns.md b/docs/04-implementation/01-known-knowns-and-unknowns.md
index 0c8e50f..084afc8 100644
--- a/docs/04-implementation/01-known-knowns-and-unknowns.md
+++ b/docs/04-implementation/01-known-knowns-and-unknowns.md
@@ -67,6 +67,18 @@ listed as an unknown-unknown has a *method* by which we will discover it.
* The client's page-table contract for server-owned ranges is explicit
(`COPY_SERVER_RESERVED_PDES`) — and its leaf levels are *empty*, i.e. ours to
populate.
+* **End-to-end CUDA works in emulation mode.** A stock CUDA program in the
+ guest (device properties, `cudaMalloc`, a GPU-authored string kernel, two
+ 4 MB host-to-device copies, a 1 M-element `vectorAdd`, and the
+ device-to-host copy-back) completes with `bad=0` against the host-computed
+ reference and the expected checksum. The host-side bridge executes the
+ client's kernels on the physical GPU, and the emulated CE path stages the
+ data in and out. Two correctness invariants are required and verified:
+ a decoded copy executes at most once per GPFIFO entry identity, and every
+ virtual operand is translated per 4 KB page (see
+ [`../03-architecture/06-host-module-and-hybrid-model.md`](../03-architecture/06-host-module-and-hybrid-model.md) §4.4).
+ The acceptance is deterministic across consecutive supervised cycles
+ (identical input/output hashes and checksum).
### 1.4 Operational
diff --git a/kernel/bar_emulation.c b/kernel/bar_emulation.c
index 685831a..e095d8b 100644
--- a/kernel/bar_emulation.c
+++ b/kernel/bar_emulation.c
@@ -269,15 +269,17 @@ int nv_vgpu_bar_emul_init(struct nv_vgpu_device *vgpu)
/*
* 0x100ce0 NV_PFB_PRI_MMU_LOCAL_MEMORY_RANGE - usable FB size.
- * Encoded as fbSize = LOWER_MAG << (LOWER_SCALE + 20). Set to
- * 256MB (matches our BAR1): mag=1, scale=8 -> 1 << 28. On real
- * hardware VBIOS devinit programs this after GFW boot; the guest
- * driver's kgspPopulateWprMeta reads it via kmemsysGetUsableFbSize
- * and uses it to compute the FRTS offset (fbSize - reservedFromEnd).
- * With it 0, the FRTS offset underflows and FWSEC verification
- * fails with "WPR2 initialized at an unexpected location".
+ * Encoded as fbSize = LOWER_MAG << (LOWER_SCALE + 20) (LOWER_SCALE
+ * bits 3:0, LOWER_MAG bits 9:4; kmemsysReadUsableFbSize_GP102).
+ * Set to 512MB (mag=1, scale=9 -> 1 << 29): on real hardware VBIOS
+ * devinit programs this after GFW boot; the guest driver's
+ * kgspPopulateWprMeta reads it via kmemsysGetUsableFbSize and uses
+ * it to compute the FRTS offset (fbSize - reservedFromEnd). With it
+ * 0, the FRTS offset underflows and FWSEC verification fails with
+ * "WPR2 initialized at an unexpected location". Must stay in
+ * lockstep with GSCI fb_length/frtsOffset and wpr2_addr_lo.
*/
- *(u32 *)((u8 *)ctx->bar0_shadow + 0x100ce0) = 0x00000018;
+ *(u32 *)((u8 *)ctx->bar0_shadow + 0x100ce0) = 0x00000019;
/*
* NV_PCFG (BAR0+0x88000) - PCI config space mirror the driver reads
diff --git a/kernel/bar_init.c b/kernel/bar_init.c
index 9926e2b..5e4b0a0 100644
--- a/kernel/bar_init.c
+++ b/kernel/bar_init.c
@@ -5,6 +5,7 @@
#include "nvidia-vgpu-vfio.h"
#include "bar_emulation.h"
+#include "gsp_phys_gpu.h"
#include "irq.h"
int nv_vgpu_bar_init(struct nv_vgpu_device *vgpu)
@@ -75,12 +76,34 @@ int nv_vgpu_bar_init(struct nv_vgpu_device *vgpu)
if (bar_addr && bar_size) {
vgpu->bars[2].gpa = bar_addr;
vgpu->bars[2].size = bar_size;
- vgpu->bars[2].vaddr = ioremap_wc(bar_addr, bar_size);
- if (!vgpu->bars[2].vaddr) {
- nv_vgpu_err(vgpu, "Failed to map BAR2");
- return -ENOMEM;
+ /*
+ * Emulation-only mode (gsp_phys_gpu=0): the host nvidia driver
+ * owns the physical GPU. Never ioremap the real BAR2 nor expose
+ * its PFNs to the guest - guest BAR2 MMIO against a device the
+ * distribution driver owns is the suspected trigger of the Xid
+ * 140 uncorrectable-ECC incident (2026-10-03). Back the region
+ * with a private host-RAM shadow instead, exactly like BAR1/BAR3.
+ * Claim/backend mode keeps the real mapping.
+ */
+ if (gsp_phys_gpu_enabled()) {
+ vgpu->bars[2].vaddr = ioremap_wc(bar_addr, bar_size);
+ if (!vgpu->bars[2].vaddr) {
+ nv_vgpu_err(vgpu, "Failed to map BAR2");
+ return -ENOMEM;
+ }
+ nv_vgpu_info(vgpu, "BAR2: gpa=0x%llx, size=0x%llx, vaddr=%p (real, claim mode)",
+ bar_addr, bar_size, vgpu->bars[2].vaddr);
+ } else {
+ vgpu->bars[2].vaddr = vzalloc(bar_size);
+ vgpu->bars[2].emulated = true;
+ if (!vgpu->bars[2].vaddr) {
+ nv_vgpu_err(vgpu, "Failed to allocate emulated BAR2 (%llu bytes)",
+ (unsigned long long)bar_size);
+ return -ENOMEM;
+ }
+ nv_vgpu_info(vgpu, "BAR2: gpa=0x%llx, size=0x%llx, emulated shadow=%p",
+ bar_addr, bar_size, vgpu->bars[2].vaddr);
}
- nv_vgpu_info(vgpu, "BAR2: gpa=0x%llx, size=0x%llx, vaddr=%p", bar_addr, bar_size, vgpu->bars[2].vaddr);
}
bar_addr = pci_resource_start(pdev, 3);
@@ -142,9 +165,17 @@ void nv_vgpu_bar_cleanup(struct nv_vgpu_device *vgpu)
vgpu->bars[3].vaddr = NULL;
}
+ /* Free the emulated BAR2 shadow (emulation-only mode) */
+ if (vgpu->bars[2].emulated && vgpu->bars[2].vaddr) {
+ vfree(vgpu->bars[2].vaddr);
+ vgpu->bars[2].vaddr = NULL;
+ }
+
for (i = 0; i < NV_VGPU_MAX_REGIONS; i++) {
if (i == 1 || i == 3)
continue; /* shadow buffers already freed above */
+ if (i == 2 && vgpu->bars[2].emulated)
+ continue; /* shadow buffer already freed above */
if (vgpu->bars[i].vaddr) {
iounmap(vgpu->bars[i].vaddr);
vgpu->bars[i].vaddr = NULL;
diff --git a/kernel/gsp_cuda_bridge.c b/kernel/gsp_cuda_bridge.c
index aaad5ee..2fac2b2 100644
--- a/kernel/gsp_cuda_bridge.c
+++ b/kernel/gsp_cuda_bridge.c
@@ -36,6 +36,7 @@
#include <linux/string.h>
#include "nvidia-vgpu-vfio.h"
+#include "gsp_cuda_bridge.h"
/* Records and limits (see tools/cuda-bridge/PROTOCOL.md, v2). */
#define CUBB2_JOB_MAGIC 0x32424243u /* 'CBB2' */
@@ -111,14 +112,7 @@ static_assert(offsetof(struct cubb2_job, prog_va) == 0x60,
static_assert(offsetof(struct cubb2_done, host_ns) == 0x70,
"cubb2_done natural layout changed");
-/* Entry points (also declared in gsp_sema_service.c). */
-bool gsp_cuda_bridge_enabled(void);
-void gsp_cuda_bridge_launch(struct nv_vgpu_device *vgpu, u32 h_object,
- const u32 qmd[64], u64 qmd_dev_va, u64 push_va,
- const u64 *recent_dsts, u32 n_recent_dsts);
-void gsp_cuda_bridge_service(struct nv_vgpu_device *vgpu);
-int gsp_cuda_bridge_init(void);
-void gsp_cuda_bridge_fini(void);
+/* Entry points are declared in gsp_cuda_bridge.h. */
/* Helpers exported by gsp_sema_service.c (declared here, not in the header,
* so the bridge stays a leaf of the sema service). */
@@ -128,6 +122,10 @@ int gsp_sema_service_va_write(struct nv_vgpu_device *vgpu, u64 va,
const void *buf, size_t len);
int gsp_sema_service_va_resolve(struct nv_vgpu_device *vgpu, u64 va, u64 *pa,
u32 *ap);
+int gsp_sema_service_va_translate(struct nv_vgpu_device *vgpu, u64 va,
+ u64 *pa, u32 *ap);
+void gsp_sema_service_va_explain(struct nv_vgpu_device *vgpu, u64 va);
+void gsp_sema_service_trace_window(u32 ms);
int gsp_sema_service_sem_read(struct nv_vgpu_device *vgpu, u64 va, u32 *val);
int gsp_sema_service_sem_write(struct nv_vgpu_device *vgpu, u32 h, u64 va,
u32 val);
@@ -148,8 +146,8 @@ struct cubb_req {
u64 push_va;
u32 qmd[64];
u32 attempts;
- u32 n_recent_dsts;
- u64 recent_dsts[CUBB2_INLINE_MAX];
+ u32 n_inl;
+ struct gsp_bridge_inline inl[GSP_BRIDGE_INLINE_MAX];
};
static struct cubb_req cubb_req[CUBB2_REQ_MAX];
@@ -190,6 +188,18 @@ struct cubb_pend {
u32 sem_flags;
struct cubb2_job job;
u8 *in; /* vmalloc'd input blob */
+ /*
+ * Deferred completion: a done record whose output write could not
+ * resolve yet (the UVM device VA's translation is learned
+ * asynchronously) is staged here and retried on later service passes
+ * for a bounded time instead of failing the job immediately.
+ */
+ bool have_done;
+ u32 done_tries;
+ unsigned long done_at;
+ u64 done_out_len;
+ u8 *done_out; /* vmalloc'd output blob */
+ struct cubb2_done done;
};
static struct cubb_pend cubb_pend[CUBB2_PEND_MAX];
@@ -215,6 +225,7 @@ static u64 cubb_dropped; /* request ring/marshal/allocation drops */
static u64 cubb_stale; /* done records for unknown seq */
static u64 cubb_malformed; /* done records with bad magic/version */
static u64 cubb_last_seq;
+static bool cubb_explained; /* one-shot walk explainer fired */
static int cubb_last_err;
static u64 cubb_launch_log; /* verbose header budget */
@@ -387,17 +398,21 @@ static int cubb_shape_kind(const u8 *p, size_t len)
/*
* Locate the launch's argument block inside the QMD parameter bank. CUDA
* writes the arguments with GRCE LOAD_INLINE_DATA at an arbitrary bank
- * offset, so never assume +0: first try the recorded recent inline-write
- * destinations that fall inside [bank_va, bank_va + bank_size), newest
- * first, then scan the bank for a mail/vadd argument shape. `want` names the
- * shape implied by grid/block (0 = unknown). On a match returns 0 and fills
- * off/kind/params/src; -ENODATA means the bank was readable but no shape was
- * found (caller may fall back to bank+0); -ENOENT means a read failed (the
- * page resolve may still be in progress) and nothing matched, so the caller
- * should retry.
+ * offset, so never assume +0: first match the payload bytes the sema service
+ * cached for the recent inline-write destinations that fall inside
+ * [bank_va, bank_va + bank_size), newest first (no VA read needed: these
+ * banks are UVM-managed VAs with no reliable translation). If the writes
+ * were split, VA-adjacent cached records are merged and shape-scanned. Then
+ * scan the bank itself for a mail/vadd argument shape as a fallback. `want`
+ * names the shape implied by grid/block (0 = unknown). On a match returns 0
+ * and fills off/kind/params/src; -ENODATA means nothing matched and no read
+ * failed (caller may fall back to bank+0); -ENOENT means a bank read failed
+ * (the page resolve may still be in progress) and nothing matched, so the
+ * caller should retry.
*/
static int cubb_find_params(struct nv_vgpu_device *vgpu, u64 bank_va,
- u32 bank_size, const u64 *dsts, u32 n_dsts,
+ u32 bank_size,
+ const struct gsp_bridge_inline *recs, u32 n_recs,
u32 want, u64 *off, u32 *kind, u8 *params,
size_t plen, const char **src)
{
@@ -412,31 +427,72 @@ static int cubb_find_params(struct nv_vgpu_device *vgpu, u64 bank_va,
if (bank_size < 16)
return -ENODATA;
- /* (a) most recently written inline destinations inside the bank */
- for (i = 0; i < n_dsts; i++) {
- u64 va = dsts[i];
+ /*
+ * (a) cached inline-write payloads, newest first. The bytes are in
+ * hand from the push decode, so a shape match needs no translation.
+ */
+ for (i = 0; i < n_recs; i++) {
+ const struct gsp_bridge_inline *r = &recs[i];
size_t need = want == 2 ? 32 : 12;
- size_t got = min_t(size_t, plen, sizeof(cand));
+ size_t got;
int k;
- if (!va || va < bank_va || va + need > hi)
- continue;
- if (va + got > hi)
- got = hi - va;
- if (gsp_sema_service_va_read(vgpu, va, cand, got)) {
- read_fail = true;
+ if (!r->dst_va || !r->len || r->dst_va < bank_va ||
+ r->dst_va + need > hi)
continue;
- }
- k = cubb_shape_kind(cand, got);
+ got = min_t(size_t, r->len, sizeof(cand));
+ if (r->dst_va + got > hi)
+ got = hi - r->dst_va;
+ k = cubb_shape_kind(r->data, got);
if (!k || (want && k != want))
continue;
- *off = va - bank_va;
+ *off = r->dst_va - bank_va;
*kind = k;
- memcpy(params, cand, min_t(size_t, got, plen));
- *src = "inline";
+ memcpy(params, r->data, min_t(size_t, got, plen));
+ *src = "cache";
return 0;
}
+ /*
+ * (a2) a large argument block can arrive as several consecutive
+ * LOAD_INLINE_DATA runs: merge VA-adjacent cached records and scan the
+ * merged bytes for a shape.
+ */
+ for (i = 0; i < n_recs; i++) {
+ u8 merged[CUBB2_PARAM_BLOB];
+ u32 mlen, j;
+ u64 base;
+ int k;
+
+ base = recs[i].dst_va;
+ if (!base || !recs[i].len || base < bank_va || base + 12 > hi)
+ continue;
+ mlen = min_t(u32, recs[i].len, sizeof(merged));
+ memcpy(merged, recs[i].data, mlen);
+ for (j = 0; j < n_recs && mlen < sizeof(merged); j++) {
+ u32 take;
+
+ if (!recs[j].len ||
+ recs[j].dst_va != base + mlen)
+ continue;
+ take = min_t(u32, recs[j].len,
+ sizeof(merged) - mlen);
+ memcpy(merged + mlen, recs[j].data, take);
+ mlen += take;
+ }
+ for (o = 0; o + 12 <= mlen; o += 4) {
+ k = cubb_shape_kind(merged + o, mlen - o);
+ if (!k || (want && k != want))
+ continue;
+ *off = base + o - bank_va;
+ *kind = k;
+ memcpy(params, merged + o,
+ min_t(u32, mlen - o, plen));
+ *src = "cache-merge";
+ return 0;
+ }
+ }
+
/* (b) scan the bank for an argument shape (straddle-safe per page) */
scan = bank_size > CUBB2_PARAM_SCAN_MAX ? CUBB2_PARAM_SCAN_MAX :
bank_size;
@@ -512,7 +568,7 @@ bool gsp_cuda_bridge_enabled(void)
*/
void gsp_cuda_bridge_launch(struct nv_vgpu_device *vgpu, u32 h_object,
const u32 qmd[64], u64 qmd_dev_va, u64 push_va,
- const u64 *recent_dsts, u32 n_recent_dsts)
+ const struct gsp_bridge_inline *recs, u32 n_recs)
{
unsigned long flags;
u64 hash = 1469598103934665603ULL;
@@ -561,13 +617,13 @@ void gsp_cuda_bridge_launch(struct nv_vgpu_device *vgpu, u32 h_object,
r->push_va = push_va;
r->attempts = 0;
memcpy(r->qmd, qmd, sizeof(r->qmd));
- r->n_recent_dsts = 0;
- if (recent_dsts && n_recent_dsts) {
- if (n_recent_dsts > CUBB2_INLINE_MAX)
- n_recent_dsts = CUBB2_INLINE_MAX;
- memcpy(r->recent_dsts, recent_dsts,
- n_recent_dsts * sizeof(u64));
- r->n_recent_dsts = n_recent_dsts;
+ r->n_inl = 0;
+ if (recs && n_recs) {
+ if (n_recs > GSP_BRIDGE_INLINE_MAX)
+ n_recs = GSP_BRIDGE_INLINE_MAX;
+ memcpy(r->inl, recs,
+ n_recs * sizeof(struct gsp_bridge_inline));
+ r->n_inl = n_recs;
}
}
cubb_req_head++;
@@ -626,9 +682,41 @@ static void cubb_log_raw_qmd(const u32 *qmd)
qmd[i], qmd[i + 1], qmd[i + 2], qmd[i + 3]);
}
+/*
+ * One-shot evidence dump when a launch's argument block cannot be located:
+ * the cached inline-write records the search had to work with (destination,
+ * length, first bytes) and the QMD's declared parameter-bank window. Decides
+ * on the next run whether the cache missed the write or the shape check was
+ * too strict.
+ */
+static u64 cubb_inline_log;
+
+static void cubb_dump_inline(const struct cubb_req *r, u64 param_va,
+ u32 param_size)
+{
+ u32 i;
+
+ if (cubb_inline_log >= 8)
+ return;
+ cubb_inline_log++;
+ pr_info("param miss: bank=0x%llx size=0x%x push=0x%llx h=0x%x records=%u\n",
+ param_va, param_size, r->push_va, r->h_object, r->n_inl);
+ for (i = 0; i < r->n_inl; i++) {
+ const struct gsp_bridge_inline *rec = &r->inl[i];
+ u32 show = min_t(u32, rec->len, 16);
+
+ pr_info("param miss: [%u] dst=0x%llx len=%u data16=%*phN\n",
+ i, rec->dst_va, rec->len, show, rec->data);
+ }
+}
+
static int cubb_marshal_one(struct nv_vgpu_device *vgpu)
{
- struct cubb_req r;
+ /*
+ * Static: the request carries 16 inline-payload records (~1.2 KB) and
+ * this runs on the single service thread only (frame-size audit).
+ */
+ static struct cubb_req r;
struct cubb2_job job;
unsigned long flags;
u32 gx, gy, gz, bx, by, bz, shared;
@@ -687,12 +775,13 @@ static int cubb_marshal_one(struct nv_vgpu_device *vgpu)
bz == 1)
hint = 2;
- rc = cubb_find_params(vgpu, param_va, param_size, r.recent_dsts,
- r.n_recent_dsts, hint, &params_off, &found_kind,
+ rc = cubb_find_params(vgpu, param_va, param_size, r.inl,
+ r.n_inl, hint, &params_off, &found_kind,
params, plen, &params_src);
if (rc) {
if (cubb_req_retry(&r, rc))
return 0;
+ cubb_dump_inline(&r, param_va, param_size);
if (rc != -ENODATA) {
cubb_req_drop_head();
cubb_failed++;
@@ -796,6 +885,36 @@ static int cubb_marshal_one(struct nv_vgpu_device *vgpu)
a, b, rc);
return 0;
}
+ /*
+ * Input-path proof: hash exactly what the bridge read from the
+ * guest's device VAs (FNV-1a over the raw bytes) plus a float
+ * NaN count per half, so the guest-side data and the daemon's
+ * received data can be compared without guessing.
+ */
+ {
+ uint64_t h = 1469598103934665603ULL;
+ u32 k, nan_a = 0, nan_b = 0;
+ const u32 *w = (const u32 *)in;
+
+ for (k = 0; k < in_bytes / 4; k++) {
+ h ^= w[k];
+ h *= 1099511628211ULL;
+ }
+ for (k = 0; k < in_bytes / 8; k++) {
+ u32 wa = w[k];
+ u32 wb = w[in_bytes / 8 + k];
+
+ if ((wa & 0x7f800000u) == 0x7f800000u &&
+ (wa & 0x007fffffu))
+ nan_a++;
+ if ((wb & 0x7f800000u) == 0x7f800000u &&
+ (wb & 0x007fffffu))
+ nan_b++;
+ }
+ pr_info("bridge-in: a=0x%llx b=0x%llx bytes=%llu hash=%016llx nan_a=%u nan_b=%u\n",
+ a, b, in_bytes, (unsigned long long)h,
+ nan_a, nan_b);
+ }
}
if (!n || !out_bytes || out_bytes > CUBB2_OUT_MAX) {
@@ -816,6 +935,10 @@ static int cubb_marshal_one(struct nv_vgpu_device *vgpu)
params_off, dst_va, n);
rc = gsp_sema_service_va_resolve(vgpu, dst_va, &dst_pa, &dst_ap);
+ if (rc && !cubb_explained) {
+ cubb_explained = true;
+ gsp_sema_service_va_explain(vgpu, dst_va);
+ }
memset(&job, 0, sizeof(job));
job.magic = CUBB2_JOB_MAGIC;
@@ -859,6 +982,12 @@ static int cubb_marshal_one(struct nv_vgpu_device *vgpu)
cubb_pend[slot].sem_flags = sem_flags;
cubb_pend[slot].job = job;
cubb_pend[slot].in = in;
+ /* A reused slot must never inherit the previous job's staged done. */
+ cubb_pend[slot].have_done = false;
+ cubb_pend[slot].done_out = NULL;
+ cubb_pend[slot].done_out_len = 0;
+ cubb_pend[slot].done_tries = 0;
+ memset(&cubb_pend[slot].done, 0, sizeof(cubb_pend[slot].done));
cubb_jobs++;
cubb_last_seq = job.seq;
mutex_unlock(&cubb_lock);
@@ -905,20 +1034,65 @@ static bool cubb_done_pop(struct cubb2_done *out)
return have;
}
-static void cubb_finish_one(struct nv_vgpu_device *vgpu,
- const struct cubb2_done *d)
+/* How long a staged done record waits for its destination translation (and a
+ * consistent read-back) before the job is declared failed. The service thread
+ * passes frequently, so a 300 s bound is tens of thousands of attempts. */
+#define CUBB2_FINISH_RETRY_MS 300000
+
+/* How long completion waits for a VA translation to appear (the guest's
+ * sysmem-backed device buffers are published asynchronously by the pin
+ * stream) before falling back to shadow-only delivery. */
+#define CUBB2_XLATE_GRACE_MS 5000
+
+/* Release a pending slot and compact the younger entries down one slot. */
+static void cubb_pend_release(struct cubb_pend *p, int idx)
+{
+ vfree(p->in);
+ p->in = NULL;
+ vfree(p->done_out);
+ p->done_out = NULL;
+ /*
+ * Clear the deferred-completion state together with the blob. Leaving
+ * have_done/done_out_len set on a freed slot is fatal: the next job
+ * staged into the same slot then looks like it already has a staged
+ * done record and the completion path copies from a NULL blob
+ * (observed: svc_vas_write NULL deref on the second launch, killing
+ * the service thread while it held cubb_lock).
+ */
+ p->have_done = false;
+ p->done_out_len = 0;
+ p->done_tries = 0;
+ p->valid = false;
+ cubb_out_len = 0;
+ if (idx >= 0) {
+ int j = idx;
+
+ while (j + 1 < CUBB2_PEND_MAX && cubb_pend[j + 1].valid) {
+ cubb_pend[j] = cubb_pend[j + 1];
+ j++;
+ }
+ cubb_pend[j].in = NULL;
+ cubb_pend[j].done_out = NULL;
+ cubb_pend[j].have_done = false;
+ cubb_pend[j].done_out_len = 0;
+ cubb_pend[j].done_tries = 0;
+ cubb_pend[j].valid = false;
+ }
+}
+
+/*
+ * Stage a daemon done record: validate it against the pending job and copy the
+ * output blob out of the shared accumulation buffer (so the next job's writes
+ * cannot overwrite it while completion retries). Called under cubb_lock.
+ */
+static void cubb_stage_done(const struct cubb2_done *d)
{
struct cubb_pend *p;
int idx = -1;
- int rc = 0;
- u32 sem_val = 0;
- u32 sem_rc = 0;
u64 wrote;
- mutex_lock(&cubb_lock);
p = cubb_pend_lookup(d->seq, &idx);
if (!p) {
- mutex_unlock(&cubb_lock);
cubb_stale++;
pr_info_ratelimited("stale done seq=%llu (no pending job)\n",
d->seq);
@@ -930,7 +1104,8 @@ static void cubb_finish_one(struct nv_vgpu_device *vgpu,
cubb_last_err = (int)d->err_code;
pr_info("failed seq=%llu err=%u status=%u op=%u\n",
d->seq, d->err_code, d->status, p->opcode);
- goto out;
+ cubb_pend_release(p, idx);
+ return;
}
wrote = min_t(u64, p->out_bytes, cubb_out_len);
@@ -939,45 +1114,134 @@ static void cubb_finish_one(struct nv_vgpu_device *vgpu,
cubb_last_err = -ENODATA;
pr_info("failed seq=%llu err=%d (daemon wrote %llu of %llu out bytes)\n",
d->seq, -ENODATA, wrote, p->out_bytes);
- goto out;
+ cubb_pend_release(p, idx);
+ return;
}
- rc = gsp_sema_service_va_write(vgpu, p->dst_va, cubb_out_buf, wrote);
- if (rc) {
+ p->done_out = vmalloc(wrote);
+ if (!p->done_out) {
cubb_failed++;
- cubb_last_err = rc;
- pr_info("failed seq=%llu err=%d (output write va=0x%llx len=%llu)\n",
- d->seq, rc, p->dst_va, wrote);
- goto out;
+ cubb_last_err = -ENOMEM;
+ pr_info("failed seq=%llu err=%d (done out alloc %llu bytes)\n",
+ d->seq, -ENOMEM, wrote);
+ cubb_pend_release(p, idx);
+ return;
}
+ memcpy(p->done_out, cubb_out_buf, wrote);
+ p->done_out_len = wrote;
/*
- * Read-back check over the whole output: proves the bytes landed
- * through the translation the guest will read (vidmem shadow or
- * guest RAM). A missing/mismatching copy must NOT be signalled
- * complete, or the guest reads stale data and runs on.
+ * Output-path proof: hash the blob the daemon produced exactly as it
+ * arrives (before any delivery translation) so the daemon's OUTCHECK
+ * hash and the guest's read-back can be compared end to end.
+ */
+ if (wrote >= (256u << 10)) {
+ uint64_t h = 1469598103934665603ULL;
+ u64 k;
+ const u32 *w = (const u32 *)p->done_out;
+
+ for (k = 0; k < wrote / 4; k++) {
+ h ^= w[k];
+ h *= 1099511628211ULL;
+ }
+ pr_info("bridge-out: seq=%llu bytes=%llu hash=%016llx\n",
+ d->seq, wrote, (unsigned long long)h);
+ }
+ p->done = *d;
+ p->have_done = true;
+ p->done_tries = 0;
+ p->done_at = jiffies;
+ cubb_out_len = 0;
+}
+
+/*
+ * Try to complete one staged job: write the output through the VA
+ * translation, read it back, then apply the release semaphore. A translation
+ * miss (-ENOENT/-EAGAIN) or a read-back mismatch (-EIO) is retried until the
+ * bounded deadline (the UVM device VA's translation is learned
+ * asynchronously); every other error fails immediately. Called under
+ * cubb_lock; returns true when the job was released.
+ */
+static bool cubb_finish_try(struct nv_vgpu_device *vgpu, struct cubb_pend *p,
+ int idx)
+{
+ const struct cubb2_done *d = &p->done;
+ u32 sem_val = 0;
+ u32 sem_rc = 0;
+ bool retryable;
+ int rc;
+
+ /*
+ * Defensive: a staged done record with output bytes but no blob is
+ * invalid state (a slot reused before this was fixed). Never copy
+ * from a NULL source - release the job instead of oopsing the
+ * service thread while it holds cubb_lock.
+ */
+ if (!p->done_out && p->done_out_len) {
+ cubb_failed++;
+ cubb_last_err = -EINVAL;
+ pr_info("seq=%llu invalid staged done (out=%llu, no blob), dropping\n",
+ d->seq, p->done_out_len);
+ cubb_pend_release(p, idx);
+ return true;
+ }
+
+ /*
+ * The guest's read-back may be a CPU read of a guest page named by the
+ * VA's translation (UVM sysmem-backed device buffers), so prefer a
+ * translated delivery. While no translation exists, stage the bytes in
+ * the VA shadow and retry for a bounded grace window (the pin stream
+ * publishes guest mappings asynchronously); after the window fall back
+ * to shadow-only delivery so a shadow-served read still works and the
+ * failure stays visible.
*/
{
+ u64 xpa = 0;
+ u32 xap = 0;
+
+ if (gsp_sema_service_va_translate(vgpu, p->dst_va, &xpa, &xap)) {
+ (void)gsp_sema_service_va_write(vgpu, p->dst_va,
+ p->done_out,
+ p->done_out_len);
+ if (time_before(jiffies,
+ p->done_at +
+ msecs_to_jiffies(CUBB2_XLATE_GRACE_MS))) {
+ p->done_tries++;
+ if ((p->done_tries & 0x7f) == 1)
+ pr_info("seq=%llu no translation yet, staged in shadow (retry=%u)\n",
+ d->seq, p->done_tries);
+ return false;
+ }
+ }
+ }
+
+ rc = gsp_sema_service_va_write(vgpu, p->dst_va, p->done_out,
+ p->done_out_len);
+ if (!rc) {
+ /*
+ * Read-back check over the whole output: proves the bytes
+ * landed through the translation the guest will read (vidmem
+ * shadow or guest RAM). A missing/mismatching copy must NOT
+ * be signalled complete, or the guest reads stale data and
+ * runs on.
+ */
u8 check[256];
size_t off = 0;
- while (off < wrote) {
- size_t clen = min_t(size_t, wrote - off,
+ while (off < p->done_out_len) {
+ size_t clen = min_t(size_t,
+ p->done_out_len - off,
sizeof(check));
if (gsp_sema_service_va_read(vgpu, p->dst_va + off,
check, clen) ||
- memcmp(check, cubb_out_buf + off, clen)) {
- cubb_failed++;
- cubb_last_err = -EIO;
- pr_info("failed seq=%llu err=%d (output read-back mismatch va=0x%llx off=%zu len=%zu)\n",
- d->seq, -EIO, p->dst_va, off, clen);
- goto out;
+ memcmp(check, p->done_out + off, clen)) {
+ rc = -EIO;
+ break;
}
off += clen;
}
}
-
- if ((p->sem_flags & 4u) && p->sem_va) {
+ if (!rc && (p->sem_flags & 4u) && p->sem_va) {
sem_val = p->sem_payload;
if (p->sem_flags & 1u) {
u32 cur = 0;
@@ -987,16 +1251,36 @@ static void cubb_finish_one(struct nv_vgpu_device *vgpu,
}
sem_rc = gsp_sema_service_sem_write(vgpu, p->h_object,
p->sem_va, sem_val);
- if (sem_rc) {
- cubb_failed++;
- cubb_last_err = sem_rc;
- pr_info("failed seq=%llu err=%d sem_va=0x%llx pay=%u\n",
- d->seq, sem_rc, p->sem_va, sem_val);
- goto out;
+ if (sem_rc)
+ rc = sem_rc;
+ }
+
+ retryable = (rc == -ENOENT || rc == -EAGAIN || rc == -EIO);
+ p->done_tries++;
+ if (rc && retryable &&
+ time_before(jiffies,
+ p->done_at + msecs_to_jiffies(CUBB2_FINISH_RETRY_MS))) {
+ /* Periodic resolve diagnostic: proves whether the
+ * translation ever becomes available while we retry. */
+ if ((p->done_tries & 0x7f) == 1) {
+ u64 dpa = 0;
+ u32 dap = 0;
+ int rrc = gsp_sema_service_va_resolve(vgpu,
+ p->dst_va,
+ &dpa, &dap);
+
+ pr_info("seq=%llu dst resolve retry=%u rc=%d pa=0x%llx ap=%u\n",
+ d->seq, p->done_tries, rrc, dpa, dap);
}
+ return false;
}
- {
+ if (rc) {
+ cubb_failed++;
+ cubb_last_err = rc;
+ pr_info("failed seq=%llu err=%d (completion va=0x%llx len=%llu tries=%u)\n",
+ d->seq, rc, p->dst_va, p->done_out_len, p->done_tries);
+ } else {
u64 dpa = 0;
u32 dap = 0;
u64 sempa = 0;
@@ -1007,40 +1291,27 @@ static void cubb_finish_one(struct nv_vgpu_device *vgpu,
(void)gsp_sema_service_va_resolve(vgpu, p->sem_va,
&sempa, &semap);
cubb_completed++;
- pr_info("completed seq=%llu op=%u dst=0x%llx dstpa=0x%llx dstap=%u out_bytes=%llu/%llu sem_va=0x%llx sempa=0x%llx pay=%u bad=%u kernel_ns=%llu\n",
- d->seq, p->opcode, p->dst_va, dpa, dap, wrote,
- p->out_bytes, p->sem_va, sempa, sem_val, d->bad,
- d->kernel_ns);
- }
-
-out:
- vfree(p->in);
- p->in = NULL;
- p->valid = false;
- cubb_out_len = 0;
- if (idx >= 0) {
- /* compact: shift the younger entries down one slot */
- int j = idx;
-
- while (j + 1 < CUBB2_PEND_MAX && cubb_pend[j + 1].valid) {
- cubb_pend[j] = cubb_pend[j + 1];
- j++;
- }
- cubb_pend[j].in = NULL;
- cubb_pend[j].valid = false;
+ pr_info("completed seq=%llu op=%u dst=0x%llx dstpa=0x%llx dstap=%u out_bytes=%llu/%llu sem_va=0x%llx sempa=0x%llx pay=%u bad=%u kernel_ns=%llu tries=%u\n",
+ d->seq, p->opcode, p->dst_va, dpa, dap,
+ p->done_out_len, p->out_bytes, p->sem_va, sempa,
+ sem_val, d->bad, d->kernel_ns, p->done_tries);
+ /* Record how the guest reads the delivered buffer back. */
+ gsp_sema_service_trace_window(10000);
}
- mutex_unlock(&cubb_lock);
+ cubb_pend_release(p, idx);
+ return true;
}
/*
* Service-thread hook: marshal at most one queued launch into a daemon-visible
- * job, then resolve any completed jobs. Called every service pass; a no-op
- * when no launch has been seen.
+ * job, stage any returned done records, then drive staged completions. Called
+ * every service pass; a no-op when no launch has been seen.
*/
void gsp_cuda_bridge_service(struct nv_vgpu_device *vgpu)
{
struct cubb2_done d;
int guard = 0;
+ int i;
if (!gsp_cuda_bridge_enabled())
return;
@@ -1052,8 +1323,18 @@ void gsp_cuda_bridge_service(struct nv_vgpu_device *vgpu)
cubb_malformed++;
continue;
}
- cubb_finish_one(vgpu, &d);
+ mutex_lock(&cubb_lock);
+ cubb_stage_done(&d);
+ mutex_unlock(&cubb_lock);
+ }
+
+ mutex_lock(&cubb_lock);
+ for (i = 0; i < CUBB2_PEND_MAX; i++) {
+ if (cubb_pend[i].valid && cubb_pend[i].have_done &&
+ cubb_finish_try(vgpu, &cubb_pend[i], i))
+ i = -1; /* compaction shifted the entries; rescan */
}
+ mutex_unlock(&cubb_lock);
}
/* ------------------------------------------------------------------ */
@@ -1319,6 +1600,8 @@ void gsp_cuda_bridge_fini(void)
for (i = 0; i < CUBB2_PEND_MAX; i++) {
vfree(cubb_pend[i].in);
cubb_pend[i].in = NULL;
+ vfree(cubb_pend[i].done_out);
+ cubb_pend[i].done_out = NULL;
cubb_pend[i].valid = false;
}
mutex_unlock(&cubb_lock);
diff --git a/kernel/gsp_cuda_bridge.h b/kernel/gsp_cuda_bridge.h
new file mode 100644
index 0000000..2fb63d7
--- /dev/null
+++ b/kernel/gsp_cuda_bridge.h
@@ -0,0 +1,34 @@
+/* SPDX-License-Identifier: GPL-2.0 */
+#ifndef _GSP_CUDA_BRIDGE_H_
+#define _GSP_CUDA_BRIDGE_H_
+
+#include <linux/types.h>
+
+struct nv_vgpu_device;
+
+/*
+ * One decoded GRCE inline-data write: the destination VA and the payload bytes
+ * carried in the push itself. The sema service records these at decode time
+ * (the bytes are already in hand); the CUDA bridge uses them to locate a
+ * launch's argument block without a VA read, which is unavailable for
+ * UVM-managed VAs (no CPU pagemap, no published page directory, no reliable
+ * targeted resolve).
+ */
+#define GSP_BRIDGE_INLINE_MAX 16
+#define GSP_BRIDGE_INLINE_BYTES 64
+
+struct gsp_bridge_inline {
+ u64 dst_va;
+ u32 len;
+ u8 data[GSP_BRIDGE_INLINE_BYTES];
+};
+
+bool gsp_cuda_bridge_enabled(void);
+void gsp_cuda_bridge_launch(struct nv_vgpu_device *vgpu, u32 h_object,
+ const u32 qmd[64], u64 qmd_dev_va, u64 push_va,
+ const struct gsp_bridge_inline *recs, u32 n_recs);
+void gsp_cuda_bridge_service(struct nv_vgpu_device *vgpu);
+int gsp_cuda_bridge_init(void);
+void gsp_cuda_bridge_fini(void);
+
+#endif /* _GSP_CUDA_BRIDGE_H_ */
diff --git a/kernel/gsp_emulation.c b/kernel/gsp_emulation.c
index 5b75ec6..b1c67a4 100644
--- a/kernel/gsp_emulation.c
+++ b/kernel/gsp_emulation.c
@@ -742,14 +742,15 @@ int gsp_emu_init(struct nv_vgpu_device *vgpu)
/* Physical GSP boot emulation state. wpr2_addr_lo matches what the
* guest driver expects WPR2_ADDR_LO to read after FWSEC FRTS:
- * frtsOffset = fbSize - reservedFromEnd. With fbSize = 256MB
- * (LOCAL_MEMORY_RANGE) and the 570.144 image's 2MB reservation,
- * frtsOffset = 0x10000000 - 0x200000 = 0xfe00000; the register's
- * _VAL field (bits31:4) must equal frtsOffset >> 12 = 0xfe00, so the
- * raw register value is 0xfe000. */
+ * frtsOffset = fbSize - reservedFromEnd. With fbSize = 512MB
+ * (LOCAL_MEMORY_RANGE / GSCI fb_length) and the 570.144 image's 2MB
+ * reservation, frtsOffset = 0x20000000 - 0x200000 = 0x1fe00000; the
+ * register's _VAL field (bits31:4) must equal frtsOffset >> 12 =
+ * 0x1fe00, so the raw register value is 0x1fe000. Kept in lockstep
+ * with GSCI fb_length/frtsOffset and NV_PFB_PRI_MMU_LOCAL_MEMORY_RANGE. */
ctx->boot_stage = 0; /* GSP_BOOT_STAGE_IDLE */
ctx->flcn_kick_count = 0;
- ctx->wpr2_addr_lo = 0x000fe000;
+ ctx->wpr2_addr_lo = 0x001fe000;
/* The guest's BAR0 is the emulation shadow (vzalloc'd by
* nv_vgpu_bar_emul_init, mapped to the guest via VFIO mmap) - NOT the
diff --git a/kernel/gsp_rm_emulation.c b/kernel/gsp_rm_emulation.c
index 1f636de..f123792 100644
--- a/kernel/gsp_rm_emulation.c
+++ b/kernel/gsp_rm_emulation.c
@@ -1054,11 +1054,12 @@ static void gsp_rm_build_static_info(u8 *buf)
/* fbRegionInfoParams (offset 344, 776 bytes) */
p = buf + TU104_GSCI_OFF_fbRegionInfoParams;
val32 = 1; memcpy(p, &val32, 4); /* numFBRegions = 1 */
- /* Region 0: base=0, limit=256MB-1, reserved=0 */
+ /* Region 0: base=0, limit=512MB-1, reserved=0 */
{
struct tu104_fb_region *region = (void *)(p + 8); /* 8-aligned */
val64 = 0; memcpy(&region->base, &val64, 8);
- val64 = 0x0FFFFFFF; memcpy(&region->limit, &val64, 8);
+ val64 = (NV_VGPU_EMUL_FB_SIZE - 1);
+ memcpy(&region->limit, &val64, 8);
val64 = 0; memcpy(&region->reserved, &val64, 8);
val32 = 0; memcpy(&region->performance, &val32, 4);
region->supportCompressed = 0;
@@ -1096,9 +1097,9 @@ static void gsp_rm_build_static_info(u8 *buf)
/* poisonFuseEnabled (offset 1216, 1 byte) */
buf[TU104_GSCI_OFF_poisonFuseEnabled] = 0;
- /* fb_length (offset 1224, 8 bytes) - 256MB (512MB broke the FWSEC
- * WPR2-location check below). */
- val64 = 0x10000000ULL;
+ /* fb_length (offset 1224, 8 bytes) - 512MB. Kept in lockstep with
+ * NV_PFB_PRI_MMU_LOCAL_MEMORY_RANGE and wpr2_addr_lo. */
+ val64 = NV_VGPU_EMUL_FB_SIZE;
memcpy(buf + TU104_GSCI_OFF_fb_length, &val64, 8);
/* fbio_mask (offset 1232, 8 bytes) - 8 FBIOs for 256-bit bus */
@@ -1166,7 +1167,7 @@ static void gsp_rm_build_static_info(u8 *buf)
/* fwWprLayoutOffset (offset 1640, 16 bytes): nonWprHeapOffset, frtsOffset */
val64 = 0;
memcpy(buf + TU104_GSCI_OFF_fwWprLayoutOffset, &val64, 8);
- val64 = 0xFE00000; /* frtsOffset = fbSize - 2MB (256MB) */
+ val64 = NV_VGPU_EMUL_FB_SIZE - 0x200000; /* frtsOffset = fbSize - 2MB */
memcpy(buf + TU104_GSCI_OFF_fwWprLayoutOffset + 8, &val64, 8);
}
@@ -1245,7 +1246,7 @@ static void gsp_rm_handle_get_static_info(struct gsp_emu_context *ctx,
nv_vgpu_info(ctx->vgpu, "vgsp-rm: GET_GSP_STATIC_INFO seq=%u -> %u bytes (fb=%lluMB, engines=0x%x/0x%x)",
sequence, TU104_GSCI_SIZE,
- 0x10000000ULL >> 20, 0x60003E02, 0x362020);
+ (unsigned long long)(NV_VGPU_EMUL_FB_SIZE >> 20), 0x60003E02, 0x362020);
if (gsp_rm_send(ctx, GSPRM_RPC_FUNC_GET_GSP_STATIC_INFO, sequence, 0,
reply, TU104_GSCI_SIZE))
@@ -1992,8 +1993,7 @@ static void gsp_rm_handle_kgr_get_context_buffers_info(
*/
static u32 tu104_fb_info_data(struct gsp_emu_context *ctx, u32 index)
{
- struct nv_vgpu_device *vgpu = ctx->vgpu;
- u64 fb = vgpu->fb_size ? vgpu->fb_size : NV_VGPU_BAR1_SIZE;
+ u64 fb = NV_VGPU_EMUL_FB_SIZE;
switch (index) {
case TU104_FB_INFO_IDX_TILE_REGION_COUNT:
diff --git a/kernel/gsp_rpc_client.c b/kernel/gsp_rpc_client.c
index 0b7f7a2..28659ff 100644
--- a/kernel/gsp_rpc_client.c
+++ b/kernel/gsp_rpc_client.c
@@ -592,7 +592,7 @@ int gsp_rpc_alloc_emul(struct gsp_client *client,
u64 va = addr_hint ? addr_hint : 0x100000000ULL; /* Default 4GB */
dma_addr_t iova;
static u64 next_vidmem_va = 0x100000; /* Start after 1MB (skip WPR2/reserved) */
- u64 fb_size = vgpu->fb_size ? vgpu->fb_size : NV_VGPU_BAR1_SIZE;
+ u64 fb_size = NV_VGPU_EMUL_FB_SIZE;
/* For vidmem allocations (golden image, channels, etc.), allocate from
* the emulated GPU's vidmem address range (0 - fb_size). The memmgr
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c
index 8b8708d..b367dcf 100644
--- a/kernel/gsp_sema_service.c
+++ b/kernel/gsp_sema_service.c
@@ -50,6 +50,7 @@
#include "gsp_rm_emulation.h"
#include "gsp_phys_booter.h"
#include "gsp_phys_gpu.h"
+#include "gsp_cuda_bridge.h"
#define SEMA_MAP_MAX 65536
#define SEMA_MAP_SIZE_BITS 16
@@ -183,11 +184,27 @@ struct svc_channel {
* apart, gating every later entry). Cleared
* when put wraps, so a recycled slot in the
* next epoch is processed normally. */
+ u64 slot_ops_ident[SVC_GPFIFO_MAX_BITS];
+ /* Raw GPFIFO entry words (e0 | e1<<32) whose
+ * operations were last executed from this
+ * slot. A rewind replay of the same entry
+ * (same words) must NOT re-run its ops: CE
+ * copies read a guest staging buffer the
+ * client refills per chunk, so re-executing
+ * a consumed copy writes the wrong chunk
+ * (observed: HtoD copies replayed after the
+ * staging buffer moved on, corrupting the
+ * vadd input). Releases still re-apply; only
+ * op execution is once per entry identity.
+ * 0 = no entry executed from this slot yet. */
u32 put_prev; /* put value at the previous pass (wrap detect) */
u64 last_rel_va; /* last tracking-sem VA seen in a readable push */
u32 logged_expected; /* last pool value logged (rate limit) */
unsigned long rescue_last; /* rate limit for the full-RAM rescue scan */
u32 rescue_fail; /* consecutive failed rescues for the entry */
+ unsigned long push_hold_until; /* absolute deadline of the bounded
+ * work-before-release hold for an
+ * unreadable CE entry (0 = none) */
u32 op_retry; /* passes an entry has been held while an operand
* targeted scan was in progress (0 = none) */
unsigned long defer_since; /* first failure of the pending GR entry
@@ -201,6 +218,30 @@ struct svc_channel {
u64 pushpg_va; /* last push page logged as pin-vs-window
* divergent (rate limit: one line per page) */
u8 pushpg_n; /* pushpg lines emitted for this channel (cap) */
+ u64 unread_va; /* entry VA inside its expensive-wait window
+ * (read + rescue + hold were already paid
+ * for it); pin-batch rewinds replay entries
+ * and must not pay that again per replay
+ * (observed: 223 replays of one 12.8 KB
+ * entry in 9 s, starving the other
+ * channels). A fresh injected/validated
+ * pin, a new put, or window expiry releases
+ * the wait. 0 = none */
+ u32 unread_idx; /* GPFIFO slot of unread_va: identity for the
+ * cheap-wait guard (the ring put advances
+ * while the entry is pending, so the put is
+ * not a stable identity) */
+ unsigned long unread_until; /* expiry of the cheap-wait window */
+ u64 unread_ipin_pa; /* injected-pin PA when the cheap-wait window
+ * opened: a replay waits cheaply while the same
+ * (possibly wrong) injection is all there is,
+ * and re-runs the rescue only when a different
+ * translation appears (0 = none seen) */
+ u32 content_max; /* highest release payload the content scans
+ * have seen for this channel's slot: the
+ * real push watermark, used to re-sync a
+ * drifted expectation instead of demanding
+ * content that does not exist yet */
u32 pending_rel_pay; /* release payload decoded while counter_pa
* was still unknown (the injected pin for the
* counter page can land a moment after the
@@ -395,6 +436,14 @@ struct sema_svc {
u64 memsets_done;
u64 memsets_nomap;
u64 memsets_fail;
+ /* VA shadow: bytes served and unresolved device writes rescued */
+ u64 vas_writes;
+ u64 vas_reads;
+ u64 vas_shadowed;
+ u64 vas_flushed; /* first bridge flush logged */
+ /* op trace window: full op/LDA logging after a bridge completion */
+ unsigned long op_trace_until;
+ u32 op_trace_log;
u64 ring_ops; /* op runs from ring-decoded pushes */
u64 lda_log; /* LAUNCH_DMA inventory budget */
u64 rescue_rej_log; /* rescue rejection budget */
@@ -407,6 +456,7 @@ struct sema_svc {
u32 op_tr_src; /* provenance of last svc_op_translate */
u64 push_pat_log; /* push-content pattern budget */
u64 dma_skip_log; /* unrecognized-DMA budget */
+ u64 xfer_raw_log; /* operandless transfer raw-dump budget */
u64 grce_log; /* GRCE inline-write budget */
u32 cur_ch; /* channel of the op being run */
u64 entry_log; /* per-push summary budget */
@@ -415,6 +465,10 @@ struct sema_svc {
u64 sem_write_reresolve; /* writes re-resolved after read-back */
u64 pushpg_fallback; /* injected pin rejected by content */
u64 gpfifo_skipped; /* stale entries consumed */
+ u64 gpfifo_rescued; /* unreadable push made readable
+ * by the content rescue */
+ u64 resc_resync; /* expectation re-synced down to
+ * the content watermark */
u64 releases_seen;
u64 releases_applied;
u64 releases_recovered; /* expected-value fallback */
@@ -565,6 +619,50 @@ static void svc_log_stalled_va(u64 va)
}
}
+/*
+ * Deferred readability probe for unreadable CE pushes that were consumed via
+ * the fake-release fallback. The pump fakes immediately (work-before-release
+ * cannot hold the channel); this queue re-attempts the read later in the
+ * service loop and logs, bounded, whether the page's content ever became
+ * resolvable - and after how long. Zero behaviour change: it only observes.
+ * The answer decides between "pump ahead of the guest's writes" (content
+ * appears moments later -> a bounded hold is viable) and "content never
+ * lands in guest RAM" (needs a different source, e.g. a writer op or the
+ * vidmem shadow).
+ */
+#define SVC_RETRY_Q 64
+#define SVC_RETRY_GIVEUP_MS 5000
+
+struct svc_retry {
+ u64 va;
+ u32 len;
+ u32 want;
+ unsigned long at;
+ unsigned long last_try;
+ u8 logged;
+};
+
+static struct svc_retry svc_retry_q[SVC_RETRY_Q];
+static u32 svc_retry_next;
+static u64 svc_retry_probes;
+
+static void svc_retry_enqueue(u64 va, u32 len, u32 want)
+{
+ u32 i;
+
+ for (i = 0; i < SVC_RETRY_Q; i++)
+ if (svc_retry_q[i].va == va && svc_retry_q[i].at)
+ return;
+ i = svc_retry_next;
+ svc_retry_next = (svc_retry_next + 1) % SVC_RETRY_Q;
+ svc_retry_q[i].va = va;
+ svc_retry_q[i].len = len;
+ svc_retry_q[i].want = want;
+ svc_retry_q[i].at = jiffies;
+ svc_retry_q[i].last_try = jiffies;
+ svc_retry_q[i].logged = 0;
+}
+
/* ------------------------------------------------------------------ */
/* guest memory helpers */
/* ------------------------------------------------------------------ */
@@ -848,6 +946,8 @@ void gsp_sema_service_vaspace_del(u32 h_vaspace)
* never maps into the buffer */
static int svc_map_lookup(u64 gpu_va, u64 *pa);
+static int svc_op_translate(struct nv_vgpu_device *vgpu, u64 va, u64 *addr,
+ u32 *aperture);
/* ------------------------------------------------------------------ */
/* client page-table walk */
@@ -1289,6 +1389,185 @@ static void svc_lat_store(u64 va, u64 addr, int ap)
l->valid = true;
}
+/*
+ * VA shadow: page-granular backing for device VAs whose VA->PA translation
+ * cannot be derived from the guest's page tables.
+ *
+ * The UVM client's device allocations live at high virtual addresses whose
+ * page tables the emulation cannot walk (the discovered directory is
+ * pool-local, the published directories miss the client VAs and the targeted
+ * rm_mem scan only finds sysmem records). Every write into such a VA used to
+ * be dropped and every read failed, so a kernel's output could never reach
+ * the guest and the guest's own HtoD/DtoH data never reached the host.
+ *
+ * The shadow makes those VAs byte-addressable without a translation: a write
+ * that cannot be translated is copied into a page record keyed by its VA, and
+ * reads are served from the record. Once a page has a record it stays
+ * authoritative for that page, so a translation learned later cannot make
+ * writer and reader disagree. The backing is plain anonymous pages
+ * (kzalloc), so unwritten bytes read as zero.
+ *
+ * Only service-thread users touch the shadow (copies, inline writes, bridge
+ * reads/writes), like the rest of the pump state.
+ */
+#define SVC_VAS_SLOTS 8192 /* 32 MiB cap, 4 KiB records */
+#define SVC_VAS_FLOOR (1ULL << 40) /* device-VA floor: 1 TiB */
+
+struct svc_vas_rec {
+ u64 va_page; /* 0 = empty */
+ u8 *data;
+};
+
+static struct svc_vas_rec svc_vas[SVC_VAS_SLOTS];
+static u32 svc_vas_n;
+static u64 svc_vas_log;
+static u64 svc_vas_warn;
+
+static u32 svc_vas_hash(u64 va_page)
+{
+ return (u32)((va_page >> SEMA_PAGE_SHIFT) * 2654435761u) &
+ (SVC_VAS_SLOTS - 1);
+}
+
+static struct svc_vas_rec *svc_vas_find(u64 va)
+{
+ u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1);
+ u32 h = svc_vas_hash(page);
+ u32 probe;
+
+ for (probe = 0; probe < SVC_VAS_SLOTS; probe++) {
+ struct svc_vas_rec *r = &svc_vas[(h + probe) & (SVC_VAS_SLOTS - 1)];
+
+ if (!r->va_page)
+ return NULL;
+ if (r->va_page == page)
+ return r;
+ }
+ return NULL;
+}
+
+static bool svc_vas_has(u64 va)
+{
+ return svc_vas_find(va) != NULL;
+}
+
+/*
+ * A VA may receive a new shadow record only when it is engine-only device
+ * memory: high in the client VA space, where the guest's own sysmem objects
+ * never live (observed sysmem VAs are below 6 GiB, device allocations above
+ * 100 TiB). Explicit sources (GRCE inline writes, the bridge's output write)
+ * shadow any VA because their target is device memory by construction.
+ */
+static bool svc_vas_can_shadow(u64 va)
+{
+ return va >= SVC_VAS_FLOOR;
+}
+
+static struct svc_vas_rec *svc_vas_get(u64 va, bool create)
+{
+ u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1);
+ u32 h = svc_vas_hash(page);
+ u32 probe;
+
+ for (probe = 0; probe < SVC_VAS_SLOTS; probe++) {
+ struct svc_vas_rec *r = &svc_vas[(h + probe) & (SVC_VAS_SLOTS - 1)];
+
+ if (r->va_page == page)
+ return r;
+ if (!r->va_page) {
+ if (!create)
+ return NULL;
+ r->data = kzalloc(SEMA_PAGE_SIZE, GFP_KERNEL);
+ if (!r->data)
+ return NULL;
+ r->va_page = page;
+ svc_vas_n++;
+ if (svc_vas_log < 16) {
+ svc_vas_log++;
+ pr_info("sema-vas: shadow page va=0x%llx (n=%u)\n",
+ page, svc_vas_n);
+ }
+ return r;
+ }
+ }
+ if (svc_vas_warn < 8) {
+ svc_vas_warn++;
+ pr_warn("sema-vas: shadow table full (%u pages), write va=0x%llx dropped\n",
+ svc_vas_n, va);
+ }
+ return NULL;
+}
+
+static int svc_vas_write(u64 va, const u8 *buf, size_t len)
+{
+ while (len) {
+ size_t chunk = min_t(size_t, len,
+ SEMA_PAGE_SIZE -
+ (va & (SEMA_PAGE_SIZE - 1)));
+ struct svc_vas_rec *r = svc_vas_get(va, true);
+
+ if (!r)
+ return -ENOMEM;
+ memcpy(r->data + (va & (SEMA_PAGE_SIZE - 1)), buf, chunk);
+ svc.vas_writes += chunk;
+ buf += chunk;
+ va += chunk;
+ len -= chunk;
+ }
+ return 0;
+}
+
+static int svc_vas_read(u64 va, u8 *buf, size_t len)
+{
+ while (len) {
+ size_t chunk = min_t(size_t, len,
+ SEMA_PAGE_SIZE -
+ (va & (SEMA_PAGE_SIZE - 1)));
+ struct svc_vas_rec *r = svc_vas_find(va);
+
+ if (!r)
+ return -ENOENT;
+ memcpy(buf, r->data + (va & (SEMA_PAGE_SIZE - 1)), chunk);
+ svc.vas_reads += chunk;
+ buf += chunk;
+ va += chunk;
+ len -= chunk;
+ }
+ return 0;
+}
+
+static int svc_vas_fill(u64 va, const u8 *pat, size_t elem, size_t total)
+{
+ u8 bounce[256];
+ size_t o;
+
+ for (o = 0; o < total; ) {
+ size_t chunk = min_t(size_t, total - o, sizeof(bounce));
+ size_t i;
+
+ for (i = 0; i + elem <= chunk; i += elem)
+ memcpy(bounce + i, pat, elem);
+ if (i < chunk)
+ memcpy(bounce + i, pat, chunk - i);
+ if (svc_vas_write(va + o, bounce, chunk))
+ return -ENOMEM;
+ o += chunk;
+ }
+ return 0;
+}
+
+static void svc_vas_free_all(void)
+{
+ u32 i;
+
+ for (i = 0; i < SVC_VAS_SLOTS; i++) {
+ kfree(svc_vas[i].data);
+ svc_vas[i].data = NULL;
+ svc_vas[i].va_page = 0;
+ }
+ svc_vas_n = 0;
+}
+
static int pin_lookup_tab(const struct svc_pin *tab, u32 slots, u32 probe_max,
u64 va, u64 *pa, int *ap, bool skip_win)
{
@@ -1658,6 +1937,27 @@ static int svc_lookup(struct nv_vgpu_device *vgpu, u64 va, u64 *addr, int *ap)
return -ENOENT;
}
+/*
+ * Resolve a VA for the bridge/IO helpers: try the CE operand path first (the
+ * bridge's operands are the same kind of client device pointers the copy data
+ * path resolves, including the incremental targeted scan and its pending
+ * semantics), then the general lookup with its vidmem window convention.
+ */
+static int svc_io_resolve(struct nv_vgpu_device *vgpu, u64 va, u64 *addr,
+ int *ap)
+{
+ u64 pa = 0;
+ u32 aper = 0;
+
+ if (!svc_op_translate(vgpu, va, &pa, &aper)) {
+ *addr = pa;
+ /* CE_APERTURE_VID is 0 (defined with the CE section below). */
+ *ap = (aper == 0) ? SVC_APERTURE_FB : SVC_APERTURE_SYS;
+ return 0;
+ }
+ return svc_lookup(vgpu, va, addr, ap);
+}
+
static int svc_resolve_va(struct nv_vgpu_device *vgpu, u64 va, u64 *addr,
u32 *aperture)
{
@@ -1855,13 +2155,29 @@ static int svc_pool_pa_resolve(struct nv_vgpu_device *vgpu, u32 h, u64 va,
int wap = SVC_APERTURE_SYS;
if (!svc_ipin_lookup(va, &pa, &a)) {
- svc_tr_clear_failure(page);
- if (src)
- *src = "pin";
- svc_pool_pa_note(h, va, pa, "pin");
- *addr = pa;
- *ap = a;
- return 0;
+ /*
+ * An injected pair is client-session evidence and normally
+ * exact, but a stale injection (the guest reallocated or moved
+ * the pool) resolves to an address that is not readable at
+ * all; keeping it first then poisoned every pool release
+ * (observed: deliveries stop, the tracker watermark freezes,
+ * and the guest never writes the next pushes' content). Prove
+ * the injected page is readable before trusting it; a failed
+ * proof falls through to the read-back-validated cache, the
+ * structural walk and the targeted rm_mem record.
+ */
+ u32 probe = 0;
+
+ if (!svc_read_ap(vgpu, pa, svc_ap_of(a), &probe,
+ sizeof(probe))) {
+ svc_tr_clear_failure(page);
+ if (src)
+ *src = "pin";
+ svc_pool_pa_note(h, va, pa, "pin");
+ *addr = pa;
+ *ap = a;
+ return 0;
+ }
}
if (!svc_pool_pa_lookup(va, &pa, &a)) {
svc_tr_clear_failure(page);
@@ -2384,11 +2700,12 @@ static void svc_rebuild(struct nv_vgpu_device *vgpu)
svc.releases_stale, svc.releases_jumped,
svc.walks_ok, svc.walks_ok + svc.walks_miss, svc.lat_hits,
svc_scan_cursor);
- pr_info("sema-svc: pages=%llu pushes=%llu pending=%llu nomap=%llu memset=%llu/%llu/%llu ringops=%llu denied=%llu false=%llu skip=%llu faked=%llu rec=%llu crel=%llu\n",
+ pr_info("sema-svc: pages=%llu pushes=%llu pending=%llu nomap=%llu memset=%llu/%llu/%llu ringops=%llu denied=%llu false=%llu skip=%llu resc=%llu faked=%llu rec=%llu crel=%llu\n",
svc.pages_decoded, svc.pushes_seen, svc.pushes_pending,
svc.pushes_nomap, svc.memsets_done, svc.memsets_nomap,
svc.memsets_fail, svc.ring_ops, svc.ops_denied,
- svc.false_trans, svc.gpfifo_skipped, svc.releases_faked,
+ svc.false_trans, svc.gpfifo_skipped, svc.gpfifo_rescued,
+ svc.releases_faked,
svc.releases_recovered, svc.counter_rel_applied);
pr_info("sema-svc: drain held=%llu starved=%llu forwarded=%llu faked=%llu rec=%llu\n",
svc.gpfifo_held, svc.gpfifo_starved, svc.releases_forwarded,
@@ -2400,6 +2717,9 @@ static void svc_rebuild(struct nv_vgpu_device *vgpu)
svc.gpfifo_entries_done, svc.gpfifo_nomap,
svc.vidmem_writes, svc_ring_n, svc_pb_n, svc_pb_passes,
svc.pin_n, svc.pin_drop, svc.ipin_n, svc.ipin_drop);
+ pr_info("sema-svc: va-shadow pages=%u/%u writes=%llu reads=%llu shadowed=%llu\n",
+ svc_vas_n, SVC_VAS_SLOTS, svc.vas_writes, svc.vas_reads,
+ svc.vas_shadowed);
/*
* P6b wave 1: guest-forwarding summary. Printed whenever any counter
* moved (not only on requests) so a run that captured no candidate is
@@ -2561,15 +2881,8 @@ struct svc_push_state {
u32 qmd[SVC_QMD_WORDS];
};
-/* Host CUDA bridge (gsp_cuda_bridge.c): launch transport + service hook. */
-void gsp_cuda_bridge_launch(struct nv_vgpu_device *vgpu, u32 h_object,
- const u32 qmd[SVC_QMD_WORDS], u64 qmd_dev_va,
- u64 push_va, const u64 *recent_dsts,
- u32 n_recent_dsts);
-bool gsp_cuda_bridge_enabled(void);
-void gsp_cuda_bridge_service(struct nv_vgpu_device *vgpu);
-int gsp_cuda_bridge_init(void);
-void gsp_cuda_bridge_fini(void);
+/* Host CUDA bridge (gsp_cuda_bridge.c): launch transport + service hook
+ * (prototypes in gsp_cuda_bridge.h; helpers exported to it below). */
/* VA/semaphore helpers exported to gsp_cuda_bridge.c. */
int gsp_sema_service_va_read(struct nv_vgpu_device *vgpu, u64 va, void *buf,
@@ -2578,6 +2891,10 @@ int gsp_sema_service_va_write(struct nv_vgpu_device *vgpu, u64 va,
const void *buf, size_t len);
int gsp_sema_service_va_resolve(struct nv_vgpu_device *vgpu, u64 va, u64 *pa,
u32 *ap);
+int gsp_sema_service_va_translate(struct nv_vgpu_device *vgpu, u64 va,
+ u64 *pa, u32 *ap);
+void gsp_sema_service_va_explain(struct nv_vgpu_device *vgpu, u64 va);
+void gsp_sema_service_trace_window(u32 ms);
int gsp_sema_service_sem_read(struct nv_vgpu_device *vgpu, u64 va, u32 *val);
int gsp_sema_service_sem_write(struct nv_vgpu_device *vgpu, u32 h, u64 va,
u32 val);
@@ -2699,6 +3016,7 @@ static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem,
u8 bounce[256];
size_t total = (size_t)elem * count;
size_t o;
+ bool shadow = false;
if (elem != 1 && elem != 2 && elem != 4 && elem != 8)
return -EINVAL;
@@ -2709,18 +3027,26 @@ static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem,
memcpy(pat, &value, elem);
if (!phys) {
- /* Virtual destination: translate through the pinned map first,
- * then the VA->PA map. */
+ /* Virtual destination: the VA shadow first (its pages stay
+ * authoritative), then the pinned map and the VA->PA map. */
u64 pa;
u32 aper;
- if (svc_op_translate(vgpu, addr, &pa, &aper))
- return -ENOENT;
- addr = pa;
- aperture = aper;
+ if (svc_vas_has(addr)) {
+ shadow = true;
+ } else if (svc_op_translate(vgpu, addr, &pa, &aper)) {
+ if (!svc_vas_can_shadow(addr))
+ return -ENOENT;
+ shadow = true;
+ svc.vas_shadowed++;
+ } else {
+ addr = pa;
+ aperture = aper;
+ }
}
- if (gsp_sema_dbg && svc_dst_watched(addr) && svc.copy_dst_log < 192) {
+ if (gsp_sema_dbg && !shadow && svc_dst_watched(addr) &&
+ svc.copy_dst_log < 192) {
u8 pb[8];
memset(pb, 0, sizeof(pb));
@@ -2731,6 +3057,9 @@ static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem,
svc.op_tr_src, elem, pb, svc.cur_ch);
}
+ if (shadow)
+ return svc_vas_fill(addr, pat, elem, total);
+
if (aperture == CE_APERTURE_VID) {
u8 *base = vgpu->bars[1].vaddr;
size_t size = vgpu->bars[1].size;
@@ -2852,6 +3181,13 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
{
u8 bounce[256];
size_t o;
+ bool dst_shadow = false;
+ bool src_shadow = false;
+ bool dst_virt = false, src_virt = false;
+ u64 dst_orig = dst, src_orig = src;
+ u32 dst_prov = 0, src_prov = 0;
+ bool bigcopy_this = false;
+ static u64 bigcopy_log;
if (!total || total > (256ULL << 20))
return -EINVAL;
@@ -2860,31 +3196,84 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
u64 pa;
u32 aper;
- if (svc_op_translate(vgpu, dst, &pa, &aper)) {
- if (gsp_sema_dbg && svc.op_tr_log < 64) {
- svc.op_tr_log++;
- pr_info("sema-tr-miss: DST va=0x%llx h=0x%x\n",
- dst, svc.cur_ch);
+ if (svc_vas_has(dst)) {
+ dst_shadow = true;
+ } else if (svc_op_translate(vgpu, dst, &pa, &aper)) {
+ if (!svc_vas_can_shadow(dst)) {
+ if (gsp_sema_dbg && svc.op_tr_log < 64) {
+ svc.op_tr_log++;
+ pr_info("sema-tr-miss: DST va=0x%llx h=0x%x\n",
+ dst, svc.cur_ch);
+ }
+ return -ENOENT;
}
- return -ENOENT;
+ dst_shadow = true;
+ svc.vas_shadowed++;
+ } else {
+ dst = pa;
+ dst_aperture = aper;
+ dst_prov = svc.op_tr_src;
+ dst_virt = true;
}
- dst = pa;
- dst_aperture = aper;
}
if (!src_phys) {
u64 pa;
u32 aper;
- if (svc_op_translate(vgpu, src, &pa, &aper)) {
+ if (svc_vas_has(src)) {
+ src_shadow = true;
+ } else if (svc_op_translate(vgpu, src, &pa, &aper)) {
if (gsp_sema_dbg && svc.op_tr_log < 64) {
svc.op_tr_log++;
pr_info("sema-tr-miss: SRC va=0x%llx h=0x%x\n",
src, svc.cur_ch);
}
return -ENOENT;
+ } else {
+ src = pa;
+ src_aperture = aper;
+ src_prov = svc.op_tr_src;
+ src_virt = true;
}
- src = pa;
- src_aperture = aper;
+ }
+
+ /*
+ * Large-copy resolution proof (bounded): for every multi-MB copy log
+ * the resolved source/destination, aperture, provenance and the first
+ * 16 bytes on each side before the copy. The vadd data path is exactly
+ * these copies; this separates a wrong resolution from a wrong copy.
+ */
+ if (total >= (1u << 20) && bigcopy_log < 24) {
+ u8 sb[16] = {0}, db[16] = {0};
+
+ bigcopy_log++;
+ bigcopy_this = true;
+ if (src_shadow) {
+ svc_vas_read(src_orig, sb, sizeof(sb));
+ } else if (src_aperture == CE_APERTURE_VID) {
+ u8 *base = vgpu->bars[1].vaddr;
+
+ if (base && src + sizeof(sb) <= vgpu->bars[1].size)
+ memcpy(sb, base + src, sizeof(sb));
+ } else {
+ svc_read(vgpu, src, sb, sizeof(sb));
+ }
+ if (dst_shadow) {
+ svc_vas_read(dst_orig, db, sizeof(db));
+ } else if (dst_aperture == CE_APERTURE_VID) {
+ u8 *base = vgpu->bars[1].vaddr;
+
+ if (base && dst + sizeof(db) <= vgpu->bars[1].size)
+ memcpy(db, base + dst, sizeof(db));
+ } else {
+ svc_read(vgpu, dst, db, sizeof(db));
+ }
+ pr_info("sema-bigcopy: dst_va=0x%llx -> pa=0x%llx ap=%u prov=%u shadow=%u | src_va=0x%llx -> pa=0x%llx ap=%u prov=%u shadow=%u | bytes=%zu\n",
+ dst_orig, dst, dst_aperture, dst_prov, dst_shadow,
+ src_orig, src, src_aperture, src_prov, src_shadow,
+ total);
+ pr_info("sema-bigcopy: src16=%16phN dst16_before=%16phN\n",
+ sb, db);
}
/*
@@ -2898,9 +3287,10 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
* conservative: one page per side, 4-aligned, 1..4096 bytes, and no
* same-page overlap within one aperture; anything doubtful is left to
* the ladder submission. Only the first eligible copy of the push is
- * kept (svc_fwd_cand is reset per decode pass).
+ * kept (svc_fwd_cand is reset per decode pass). Shadowed operands have
+ * no engine-reachable address, so they are never forwarding candidates.
*/
- if (!svc_fwd_cand.valid &&
+ if (!svc_fwd_cand.valid && !src_shadow && !dst_shadow &&
total >= 1 && total <= SEMA_PAGE_SIZE && !(total & 3)) {
u64 src_page = src & ~(u64)(SEMA_PAGE_SIZE - 1);
u64 dst_page = dst & ~(u64)(SEMA_PAGE_SIZE - 1);
@@ -2967,7 +3357,7 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
* source or destination starts with that 8-byte value is logged too,
* whatever its size - this finds the push carrying a known payload.
*/
- if (gsp_sema_dbg &&
+ if (gsp_sema_dbg && !src_shadow && !dst_shadow &&
((gsp_sema_copy_watch && total == gsp_sema_copy_watch) ||
gsp_sema_copy_pat)) {
u8 sb[8] = {0}, db[8] = {0};
@@ -3026,7 +3416,8 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
* answered for the source (1 pin, 2 map, 3 vaspaces, 4 discovered,
* 5 targeted RAM resolve).
*/
- if (gsp_sema_dbg && svc_dst_watched(dst) && svc.copy_dst_log < 192) {
+ if (gsp_sema_dbg && !src_shadow && !dst_shadow &&
+ svc_dst_watched(dst) && svc.copy_dst_log < 192) {
u8 peek[32];
u32 n = total < sizeof(peek) ? (u32)total : sizeof(peek);
@@ -3051,8 +3442,38 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
if (chunk > sizeof(bounce))
chunk = sizeof(bounce);
+ /*
+ * A virtual operand must be translated per 4 KB page: its
+ * pages are independent guest allocations (the pins show
+ * split runs), so a linear PA walk past the first page reads
+ * the wrong memory. Limit each bounce to the current page.
+ */
+ if (src_virt)
+ chunk = min_t(size_t, chunk,
+ 0x1000 - ((src_orig + o) & 0xfff));
+ if (dst_virt)
+ chunk = min_t(size_t, chunk,
+ 0x1000 - ((dst_orig + o) & 0xfff));
+
+ if (src_shadow) {
+ if (svc_vas_read(src_orig + o, bounce, chunk))
+ return -EIO;
+ } else if (src_virt) {
+ u64 pa = 0;
+ u32 aper = 0;
- if (src_aperture == CE_APERTURE_VID) {
+ if (svc_op_translate(vgpu, src_orig + o, &pa, &aper))
+ return -ENOENT;
+ if (aper == CE_APERTURE_VID) {
+ u8 *base = vgpu->bars[1].vaddr;
+
+ if (!base || pa + chunk > vgpu->bars[1].size)
+ return -EINVAL;
+ memcpy(bounce, base + pa, chunk);
+ } else if (svc_read(vgpu, pa, bounce, chunk)) {
+ return -EIO;
+ }
+ } else if (src_aperture == CE_APERTURE_VID) {
u8 *base = vgpu->bars[1].vaddr;
size_t size = vgpu->bars[1].size;
@@ -3064,7 +3485,25 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
return -EIO;
}
- if (dst_aperture == CE_APERTURE_VID) {
+ if (dst_shadow) {
+ if (svc_vas_write(dst_orig + o, bounce, chunk))
+ return -ENOMEM;
+ } else if (dst_virt) {
+ u64 pa = 0;
+ u32 aper = 0;
+
+ if (svc_op_translate(vgpu, dst_orig + o, &pa, &aper))
+ return -ENOENT;
+ if (aper == CE_APERTURE_VID) {
+ u8 *base = vgpu->bars[1].vaddr;
+
+ if (!base || pa + chunk > vgpu->bars[1].size)
+ return -EINVAL;
+ memcpy(base + pa, bounce, chunk);
+ } else if (svc_write(vgpu, pa, bounce, chunk)) {
+ return -EIO;
+ }
+ } else if (dst_aperture == CE_APERTURE_VID) {
u8 *base = vgpu->bars[1].vaddr;
size_t size = vgpu->bars[1].size;
@@ -3079,7 +3518,30 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
o += chunk;
}
- if (gsp_sema_dbg && gsp_sema_copy_watch &&
+ /*
+ * Large-copy landing proof: read the first 16 destination bytes back
+ * through the resolved destination so the log shows whether the copy
+ * landed where the resolver claims it should.
+ */
+ if (bigcopy_this) {
+ u8 da[16] = {0};
+
+ if (dst_shadow) {
+ svc_vas_read(dst_orig, da, sizeof(da));
+ } else if (dst_aperture == CE_APERTURE_VID) {
+ u8 *base = vgpu->bars[1].vaddr;
+
+ if (base && dst + sizeof(da) <= vgpu->bars[1].size)
+ memcpy(da, base + dst, sizeof(da));
+ } else {
+ svc_read(vgpu, dst, da, sizeof(da));
+ }
+ pr_info("sema-bigcopy: dst16_after=%16phN (dst_va=0x%llx pa=0x%llx)\n",
+ da, dst_orig, dst);
+ }
+
+ if (gsp_sema_dbg && !src_shadow && !dst_shadow &&
+ gsp_sema_copy_watch &&
total == gsp_sema_copy_watch &&
svc.copy_watch_log <= 64) {
u8 da[8] = {0};
@@ -3096,7 +3558,8 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total,
da[0], da[1], da[2], da[3], da[4], da[5], da[6],
da[7], dst);
}
- if (gsp_sema_dbg && svc_dst_watched(dst) && svc.copy_dst_log <= 192) {
+ if (gsp_sema_dbg && !src_shadow && !dst_shadow &&
+ svc_dst_watched(dst) && svc.copy_dst_log <= 192) {
u8 da[16] = {0};
u32 n = total < sizeof(da) ? (u32)total : sizeof(da);
@@ -3128,7 +3591,30 @@ static int svc_grce_inline_write(struct nv_vgpu_device *vgpu, u64 dst_va,
if (!bytes || bytes > (1u << 20))
return -EINVAL;
+ /*
+ * The shadow is authoritative for pages it already holds; otherwise
+ * resolve through the trusted sources. An unresolved destination is
+ * device memory the client written through the engine: preserve the
+ * bytes in the VA shadow instead of dropping the write (the launch
+ * parameter bank and other UVM client buffers land here).
+ */
+ if (svc_vas_has(dst_va)) {
+ if (svc_vas_write(dst_va, data, bytes))
+ return -ENOMEM;
+ svc.memsets_done++;
+ return 0;
+ }
if (svc_lookup_trusted(vgpu, dst_va & ~0xfffULL, &pa, &ap)) {
+ if (!svc_vas_write(dst_va, data, bytes)) {
+ svc.vas_shadowed++;
+ if (svc.grce_log < 16) {
+ svc.grce_log++;
+ pr_info("sema-grce: inline write dst=0x%llx bytes=%u shadowed (unresolved) h=0x%x\n",
+ dst_va, bytes, ch_h);
+ }
+ svc.memsets_done++;
+ return 0;
+ }
if (gsp_sema_dbg && svc.grce_log < 16) {
svc.grce_log++;
pr_info("sema-grce: inline write dst=0x%llx bytes=%u UNRESOLVED h=0x%x\n",
@@ -3172,6 +3658,7 @@ static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops,
u32 n_ops, u32 ch_h)
{
u32 i;
+ static u64 bigop_log;
svc.cur_ch = ch_h;
for (i = 0; i < n_ops; i++) {
@@ -3187,6 +3674,21 @@ static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops,
m->count, m->phys, m->aperture);
/*
+ * Large-copy proof (unconditional, bounded): the multi-MB
+ * HtoD/DtoH transfers are the vadd data path; log each one's
+ * operands, aperture and result so a missing or partially
+ * applied copy is visible without relying on the bounded
+ * op-trace budget.
+ */
+ if (m->bytes >= (256u << 10) && bigop_log < 48) {
+ bigop_log++;
+ pr_info("sema-bigop: type=%u dst=0x%llx src=0x%llx bytes=%llu phys=%u ap=%u srcphys=%u srcap=%u rc=%d h=0x%x\n",
+ m->type, m->dst, m->src, m->bytes, m->phys,
+ m->aperture, m->src_phys, m->src_aperture,
+ rc, ch_h);
+ }
+
+ /*
* P6b wave 2: a withheld software copy (hardware-primary) is
* not software work: it is neither a success nor a failure, so
* no memsets counter moves and no op log fires for it.
@@ -3215,6 +3717,28 @@ static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops,
m->bytes, m->phys, m->aperture, m->src,
m->src_phys, m->src_aperture, rc,
svc.cur_ch);
+ /*
+ * Diagnostic op trace: after a bridge completion (temporal
+ * window) or whenever an operand is a high device VA, log the
+ * full op so the guest's read-back encoding is on record.
+ * Bounded; no effect on execution.
+ */
+ /*
+ * Bounded; no effect on execution. A 39-byte op is the mail
+ * read-back signature: always keep room for it past the
+ * general trace budget so the DtoH's src/dst/rc is on record.
+ */
+ if ((svc.op_trace_log < 160 &&
+ (time_before(jiffies, svc.op_trace_until) ||
+ m->dst >= SVC_VAS_FLOOR ||
+ (m->type == 1 && !m->src_phys && m->src >= SVC_VAS_FLOOR))) ||
+ (m->bytes == 39 && svc.op_trace_log < 4096)) {
+ svc.op_trace_log++;
+ pr_info("sema-optrace: type=%u dst=0x%llx phys=%u ap=%u bytes=%llu src=0x%llx srcphys=%u srcap=%u rc=%d h=0x%x\n",
+ m->type, m->dst, m->phys, m->aperture,
+ m->bytes, m->src, m->src_phys, m->src_aperture,
+ rc, ch_h);
+ }
}
}
@@ -3336,11 +3860,20 @@ static u32 svc_rel_count;
static u32 svc_dbg_streams;
/*
- * Ring of recent GRCE inline-write destinations (newest at next-1), consumed
- * by the CUDA bridge to locate a launch's argument block inside the QMD
- * parameter bank. Written and read on the service thread only.
+ * Ring of recent GRCE inline-write destinations and their payload bytes
+ * (newest at next-1), consumed by the CUDA bridge to locate a launch's
+ * argument block inside the QMD parameter bank. The bytes are cached at
+ * decode time because the destination VAs are UVM-managed and often cannot
+ * be read back (no CPU pagemap, no published page directory, unreliable
+ * targeted resolve). Written and read on the service thread only.
*/
-static u64 svc_grce_dst_ring[SVC_GRCE_DST_MAX];
+struct svc_grce_rec {
+ u64 dst_va;
+ u32 len;
+ u8 data[GSP_BRIDGE_INLINE_BYTES];
+};
+
+static struct svc_grce_rec svc_grce_dst_ring[SVC_GRCE_DST_MAX];
static u32 svc_grce_dst_ring_next;
static void svc_counter_learn(struct nv_vgpu_device *vgpu, u32 h, u64 sem_va);
@@ -3630,6 +4163,33 @@ static void svc_counter_autoreg(struct nv_vgpu_device *vgpu,
derived = true;
}
ch->counter_pa = pa & ~0x3ULL;
+ /*
+ * Persist the registration into the authoritative channel array
+ * immediately, under the lock, instead of relying on the pump's
+ * end-of-pass write-back: the pump works on a snapshot and several
+ * early `continue`s (USERD not yet resolvable, ring not yet
+ * readable) skip that write-back. When that happened every pass
+ * re-registered the same slot - the "counter auto" line flooded the
+ * log ring (1.17M lines, noclaim27) and, worse, the publish path
+ * never saw a registered counter (`pub_calls=0`), so the client's
+ * tracker semaphores were never advanced and cudaMalloc hung. The
+ * authoritative store makes the registration durable and the log
+ * one-shot.
+ */
+ {
+ unsigned long flags;
+ int j;
+
+ spin_lock_irqsave(&svc.chan_lock, flags);
+ for (j = 0; j < SVC_MAX_CHANNELS; j++) {
+ if (svc.chan[j].valid &&
+ svc.chan[j].h_object == ch->h_object) {
+ svc.chan[j].counter_pa = ch->counter_pa;
+ break;
+ }
+ }
+ spin_unlock_irqrestore(&svc.chan_lock, flags);
+ }
if (derived)
pr_info("sema-chan: counter derived h=0x%x rank=%u pa=0x%llx (slot va=0x%llx)\n",
ch->h_object, ch->gr_index, ch->counter_pa, va);
@@ -3904,27 +4464,39 @@ static void svc_qmd_accum(struct nv_vgpu_device *vgpu, struct svc_channel *ch,
if (st->qmd_have == 3)
st->qmd_dev_va = (((u64)st->qmd_hi << 32) | st->qmd_lo) << 8;
if (st->qmd_n == SVC_QMD_WORDS) {
- u64 dsts[SVC_GRCE_DST_MAX];
+ /*
+ * Static: a 16-entry record array is ~1.2 KB and this is
+ * called from the single service thread only (same reason as
+ * the ops array in svc_decode_stream).
+ */
+ static struct gsp_bridge_inline recs[GSP_BRIDGE_INLINE_MAX];
u32 have, i, n = 0;
/*
- * Hand the bridge the recent inline-write destinations,
- * newest first, so it can find the argument block CUDA
- * placed inside the QMD parameter bank (never assume
- * bank+0). The current push's writes are already in the
- * ring because the QMD run follows them.
+ * Hand the bridge the recent inline-write destinations and
+ * their cached payload bytes, newest first, so it can find
+ * the argument block CUDA placed inside the QMD parameter
+ * bank (never assume bank+0) without a VA read. The current
+ * push's writes are already in the ring because the QMD run
+ * follows them.
*/
have = svc_grce_dst_ring_next < SVC_GRCE_DST_MAX ?
svc_grce_dst_ring_next : SVC_GRCE_DST_MAX;
+ if (have > GSP_BRIDGE_INLINE_MAX)
+ have = GSP_BRIDGE_INLINE_MAX;
for (i = 0; i < have; i++) {
u32 idx = (svc_grce_dst_ring_next - 1 - i) %
SVC_GRCE_DST_MAX;
+ const struct svc_grce_rec *rec = &svc_grce_dst_ring[idx];
- dsts[n++] = svc_grce_dst_ring[idx];
+ recs[n].dst_va = rec->dst_va;
+ recs[n].len = rec->len;
+ memcpy(recs[n].data, rec->data, rec->len);
+ n++;
}
gsp_cuda_bridge_launch(vgpu, ch->h_object, st->qmd,
st->qmd_dev_va, st->push_va,
- dsts, n);
+ recs, n);
st->qmd_n = 0;
st->qmd_bits = 0;
st->qmd_have = 0;
@@ -3977,6 +4549,14 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf,
svc_fwd_rel_n = 0;
svc_fwd_defer_armed = false;
svc.pages_decoded++;
+ /* Diagnostic: name every push decoded in the post-completion window. */
+ if (svc.op_trace_log < 160 &&
+ time_before(jiffies, svc.op_trace_until)) {
+ svc.op_trace_log++;
+ pr_info("sema-pushtrace: h=0x%x eng=%u va=0x%llx len=%u allow=%d trusted=%d\n",
+ ch->h_object, ch->engine_type, push_va, len,
+ allow_ops, trusted);
+ }
if (gsp_sema_dbg && ch->engine_type && svc_dbg_streams < 24) {
svc_dbg_streams++;
pr_info("sema-dbg: stream h=0x%x eng=%u va=0x%llx len=%u allow=%d w=%08x %08x %08x %08x %08x %08x %08x %08x\n",
@@ -4049,16 +4629,22 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf,
st.line_length, ch->h_object);
}
/*
- * Record the destination for the CUDA bridge's
- * argument-block search, whether or not the write was
- * executed: the launch QMD later in this push needs it
- * to find where CUDA placed the kernel arguments in
- * the QMD parameter bank.
+ * Record the destination and the payload bytes for
+ * the CUDA bridge's argument-block search, whether or
+ * not the write was executed: the launch QMD later in
+ * this push needs them to find where CUDA placed the
+ * kernel arguments in the QMD parameter bank, and the
+ * destination VA may not be readable later.
*/
if (st.have_out && st.offset_out) {
- svc_grce_dst_ring[svc_grce_dst_ring_next %
- SVC_GRCE_DST_MAX] =
- st.offset_out;
+ struct svc_grce_rec *rec =
+ &svc_grce_dst_ring[svc_grce_dst_ring_next %
+ SVC_GRCE_DST_MAX];
+
+ rec->dst_va = st.offset_out;
+ rec->len = min_t(u32, count * 4,
+ GSP_BRIDGE_INLINE_BYTES);
+ memcpy(rec->data, buf + off + 4, rec->len);
svc_grce_dst_ring_next++;
}
off += 4 + count * 4;
@@ -4083,11 +4669,32 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf,
* the inline-QMD method run for the host CUDA bridge. The run
* is skipped whole exactly as before (release/decode behaviour
* is untouched); a completed 64-word QMD queues a daemon job.
+ *
+ * The QMD run's address window [0x318, 0x420) OVERLAPS the
+ * C5B5 copy-class operand methods OFFSET_OUT_UPPER/LOWER
+ * (0x408/0x40c), LINE_LENGTH_IN (0x418) and LINE_COUNT
+ * (0x41c): CUDA binds the copy class on subchannel 4 of the
+ * same compute channel for its small copies (observed:
+ * `c0810120` SET_REMAP_CONST_A, `02810220` OFFSET_OUT count 2,
+ * `06810120` LINE_LENGTH_IN 0x2000, `c0800120` LAUNCH_DMA
+ * 0x586, `c0800120` SET_SEMAPHORE_A on the counter slot). The
+ * old unconditional `subch >= 1` intercept swallowed those
+ * operands as QMD words: the LAUNCH_DMA created no copy op,
+ * the copy never ran, yet the SET_REPORT_SEMAPHORE release
+ * still applied, so the guest read stale/zero data while its
+ * wait cleared (the invisible DtoH read-back).
+ *
+ * Only accumulate once the run has declared itself a QMD
+ * submission by carrying the ADDRESS_A/B methods; otherwise a
+ * method in this window belongs to another class's space and
+ * must fall through to the CE decoder.
*/
if (ch->engine_type == 0 && subch >= 1 &&
gsp_cuda_bridge_enabled() &&
addr >= SVC_QMD_ADDR_A &&
- addr <= SVC_QMD_DATA + (SVC_QMD_WORDS - 1) * 4) {
+ addr <= SVC_QMD_DATA + (SVC_QMD_WORDS - 1) * 4 &&
+ (st.qmd_have ||
+ (addr <= SVC_QMD_ADDR_B && addr + count * 4 > SVC_QMD_ADDR_A))) {
svc_qmd_accum(vgpu, ch, &st, addr, count,
buf + off + 4);
off += 4 + count * 4;
@@ -4411,6 +5018,27 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf,
(st.offset_out & ~0xffffffffULL) | v0;
st.have_out = 1;
break;
+ case GRCE_M_LAUNCH_DMA:
+ /*
+ * Compute-class DMA launch: the other possible
+ * encoding for a copy issued on a GR channel.
+ * Dump it (bounded), with the accumulated
+ * operand state, so one run covers both the
+ * C5B5 and the GRCE copy shapes.
+ */
+ if (ch->h_object && gsp_sema_dbg &&
+ svc.xfer_raw_log < 48) {
+ u32 dl = len < 192 ? len : 192;
+
+ svc.xfer_raw_log++;
+ pr_info("sema-xfer-raw: grce h=0x%x push=0x%llx len=%u off=0x%x v0=0x%x v1=0x%x out=0x%llx line=%u cnt=%u have_out=%u %*phN\n",
+ ch->h_object, push_va, len,
+ off, v0, v1, st.offset_out,
+ st.line_length,
+ st.gr_line_count,
+ st.have_out, dl, buf);
+ }
+ break;
default:
break;
}
@@ -4477,6 +5105,44 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf,
st.line_length, st.have_out, st.have_in,
st.have_remap, st.have_sem, n_ops);
}
+ /* Diagnostic: full operand state in the post-completion
+ * window (or for high device VAs), bounded. */
+ if (svc.op_trace_log < 160 && ch->h_object &&
+ (time_before(jiffies, svc.op_trace_until) ||
+ st.offset_out >= SVC_VAS_FLOOR ||
+ st.offset_in >= SVC_VAS_FLOOR)) {
+ svc.op_trace_log++;
+ pr_info("sema-ldatrace: h=0x%x eng=%u push=0x%llx inj=%u allow=%u dma=0x%08x xfer=%u dstphys=%u srcphys=%u smode=%u dmode=%u out=0x%llx in=0x%llx line=%u have_out=%u have_in=%u have_remap=%u have_sem=%u ops=%u\n",
+ ch->h_object, ch->engine_type, push_va,
+ trusted, allow_ops, dma, transfer,
+ DMA_DST_TYPE_PHYSICAL(dma),
+ DMA_SRC_TYPE_PHYSICAL(dma),
+ st.have_src_mode, st.have_dst_mode,
+ st.offset_out, st.offset_in,
+ st.line_length, st.have_out, st.have_in,
+ st.have_remap, st.have_sem, n_ops);
+ }
+
+ /*
+ * A transfer-shaped launch with no decoded operands is
+ * invisible today: the copy never runs, yet the push's
+ * release still applies, so the guest reads stale data
+ * while its wait clears. Dump the whole push once for a
+ * bounded number of such launches so the operand
+ * encoding (or its absence) is on record.
+ */
+ if (transfer != 0 && ch->h_object && gsp_sema_dbg &&
+ svc.xfer_raw_log < 48 &&
+ (!st.have_out || !st.line_length)) {
+ u32 dl = len < 192 ? len : 192;
+
+ svc.xfer_raw_log++;
+ pr_info("sema-xfer-raw: h=0x%x eng=%u push=0x%llx len=%u off=0x%x dma=0x%08x line=%u out=0x%llx in=0x%llx have_out=%u have_in=%u %*phN\n",
+ ch->h_object, ch->engine_type, push_va,
+ len, off, dma, st.line_length,
+ st.offset_out, st.offset_in,
+ st.have_out, st.have_in, dl, buf);
+ }
if (allow_ops && transfer != 0 && st.line_length &&
st.have_out) {
@@ -5328,6 +5994,96 @@ static int svc_read_va_trusted(struct nv_vgpu_device *vgpu, u64 va, void *buf,
}
/*
+ * Probe the deferred-retry queue: re-attempt a trusted read of each queued
+ * push VA about every 100 ms and log, once per entry, whether the content
+ * became resolvable (with age and first word) or never did within the giveup
+ * window. Observation only - no pins, no releases, no execution.
+ */
+static void svc_retry_poll(struct nv_vgpu_device *vgpu)
+{
+ u8 tmp[64];
+ u32 i;
+
+ for (i = 0; i < SVC_RETRY_Q; i++) {
+ struct svc_retry *r = &svc_retry_q[i];
+ unsigned long age;
+ int rc;
+ u32 w0 = 0;
+
+ if (!r->at)
+ continue;
+ age = jiffies_to_msecs(jiffies - r->at);
+ if (age > SVC_RETRY_GIVEUP_MS) {
+ if (!r->logged) {
+ r->logged = 1;
+ pr_info("sema-retry: va=0x%llx len=%u want=%u gave up after %lu ms (never resolvable)\n",
+ r->va, r->len, r->want, age);
+ }
+ r->at = 0;
+ continue;
+ }
+ if (time_before(jiffies, r->last_try + msecs_to_jiffies(500)))
+ continue;
+ r->last_try = jiffies;
+ svc_retry_probes++;
+ rc = svc_read_va_trusted(vgpu, r->va, tmp,
+ min_t(size_t, r->len, sizeof(tmp)));
+ if (rc && !r->logged) {
+ r->logged = 1;
+ pr_info("sema-retry: va=0x%llx len=%u want=%u GAVE UP after %lu ms (never resolvable)\n",
+ r->va, r->len, r->want, age);
+ r->at = 0;
+ continue;
+ }
+ if (rc)
+ continue;
+ if (!r->logged) {
+ u32 relpay = 0;
+ bool have_rel = false;
+
+ memcpy(&w0, tmp, 4);
+ /*
+ * Extract the release payload from the page's tail so
+ * the expectation (`want`) and the content's own
+ * payload can be compared directly: an exact-match
+ * rescue rejects the correct page when the two differ
+ * (expectation drift), while a matching payload points
+ * at the stream/parse or a different-page problem.
+ */
+ if (r->len >= 32) {
+ u8 tail[64];
+ size_t tlen = min_t(size_t, r->len,
+ sizeof(tail));
+ u32 j;
+
+ if (!svc_read_va_trusted(vgpu,
+ r->va + r->len - tlen,
+ tail, tlen)) {
+ for (j = 0; j + 16 <= tlen; j += 4) {
+ u32 hdr, hi;
+
+ memcpy(&hdr, tail + j, 4);
+ if (hdr != 0x20038090u)
+ continue;
+ memcpy(&hi, tail + j + 4, 4);
+ if (hi != 1)
+ continue;
+ memcpy(&relpay, tail + j + 12,
+ 4);
+ have_rel = true;
+ }
+ }
+ }
+ r->logged = 1;
+ pr_info("sema-retry: va=0x%llx len=%u want=%u resolvable after %lu ms w0=0x%08x rel=%d relpay=%u\n",
+ r->va, r->len, r->want, age, w0,
+ have_rel ? 1 : 0, relpay);
+ r->at = 0;
+ }
+ }
+}
+
+/*
* Linear fallback for a push-buffer window page.
*
* svc_pb_pair() pairs the window base VA with the base PA it found by content
@@ -5858,6 +6614,17 @@ u64 gsp_sema_copy_dst2;
*/
unsigned gsp_sema_ce_grace_ms = 8000;
+/*
+ * Bounded work-before-release hold (ms) for an unreadable CE entry. The
+ * deferred retry probe shows the translation/pins for a faked push typically
+ * land 0.5-2 s later, so hold the entry for a fixed absolute grace: a late
+ * translation then executes the work before its release reaches the client
+ * instead of the copy being skipped. Per-entry absolute deadline, never
+ * reset by replay/rescue paths (the noclaim9 livelock lesson); when it
+ * passes the entry is consumed exactly as before. 0 disables the hold.
+ */
+unsigned gsp_sema_push_hold_ms = 2500;
+
static unsigned long svc_gr_ring_tried; /* ranks that got a diagnostic line */
static void svc_gr_ring_autopin(struct nv_vgpu_device *vgpu,
@@ -6101,8 +6868,8 @@ static void svc_pb_pair(struct nv_vgpu_device *vgpu, struct svc_channel *ch)
* updated candidate count.
*/
static u32 svc_rescue_match_buf(const u8 *buf, u32 len, u64 base_pa, u32 slot,
- u32 want_payload, u64 entry_len, u64 *cand,
- u32 ncand)
+ u32 want_payload, u64 entry_va, u64 entry_len,
+ u64 *cand, u64 *cand_va, u32 ncand, u32 *maxp)
{
u32 o;
@@ -6123,6 +6890,14 @@ static u32 svc_rescue_match_buf(const u8 *buf, u32 len, u64 base_pa, u32 slot,
if (!payload || payload >= (1u << 24))
continue;
/*
+ * Record the highest payload this channel's slot shows in
+ * RAM (regardless of whether it is this entry's release):
+ * the real push watermark, used to re-sync an expectation
+ * that fakes have pushed ahead of the content.
+ */
+ if (maxp && payload > *maxp)
+ *maxp = payload;
+ /*
* When the channel's next expected release payload is known
* (learned from the pool), require it exactly: a recycled or
* stale copy of an older push for the same channel rank has a
@@ -6130,8 +6905,26 @@ static u32 svc_rescue_match_buf(const u8 *buf, u32 len, u64 base_pa, u32 slot,
* (observed: the second push of a channel re-pinned to the
* page of its first push, so the real push was never read).
*/
- if (want_payload && payload != want_payload)
+ if (want_payload && payload != want_payload) {
+ /*
+ * Near-miss observability: a group with this channel's
+ * slot but a different payload means content for the
+ * rank exists in RAM; the distance from `want` tells
+ * whether the expectation drifted (payloads ahead of
+ * want were faked) or the content is stale/other.
+ * Bounded to 16 lines per boot.
+ */
+ static u32 near_log;
+
+ if (gsp_sema_dbg && near_log < 32 &&
+ time_before(jiffies, svc.op_trace_until)) {
+ near_log++;
+ pr_info("sema-rescue: payload near-miss slot=%u want=%u got=%u (delta=%d)\n",
+ slot, want_payload, payload,
+ (int)payload - (int)want_payload);
+ }
continue;
+ }
memcpy(&w5, buf + o + 16, 4);
memcpy(&v5, buf + o + 20, 4);
if ((w5 >> 29) != 1 || ((w5 >> 13) & 0x7) != 4 ||
@@ -6139,14 +6932,43 @@ static u32 svc_rescue_match_buf(const u8 *buf, u32 len, u64 base_pa, u32 slot,
(v5 & 0x3) != 0)
continue;
+ /*
+ * The mapping is page-granular: the physical page holding the
+ * release group backs the VA page holding the group
+ * (entry_va + entry_len - 24), and every earlier candidate is
+ * the same page distance back in VA. Without this, a push
+ * crossing into a non-contiguous page was validated against
+ * the entry's *first* VA page and always rejected (observed:
+ * continuation page of 0x120000f58 really at 0x1033c9000,
+ * candidates pinned at the entry page instead).
+ */
rel_pa = base_pa + o;
- cand[ncand++] = rel_pa & ~(u64)(SEMA_PAGE_SIZE - 1);
- for (d = entry_len >= 24 ? entry_len - 24 : 0;
- d > 0 && ncand < SVC_MBUF_MAX; d -= SEMA_PAGE_SIZE) {
- u64 s = (rel_pa - d) & ~(u64)(SEMA_PAGE_SIZE - 1);
-
- if (s != cand[ncand - 1])
- cand[ncand++] = s;
+ {
+ u64 group_pa = rel_pa & ~(u64)(SEMA_PAGE_SIZE - 1);
+ u64 group_va = entry_va + (entry_len >= 24 ?
+ entry_len - 24 : 0);
+
+ cand[ncand] = group_pa;
+ if (cand_va)
+ cand_va[ncand] = group_va &
+ ~(u64)(SEMA_PAGE_SIZE - 1);
+ ncand++;
+ for (d = entry_len >= 24 ? entry_len - 24 : 0;
+ d > 0 && ncand < SVC_MBUF_MAX;
+ d -= SEMA_PAGE_SIZE) {
+ u64 s = (rel_pa - d) &
+ ~(u64)(SEMA_PAGE_SIZE - 1);
+
+ if (s == cand[ncand - 1])
+ continue;
+ cand[ncand] = s;
+ if (cand_va)
+ cand_va[ncand] =
+ (group_va &
+ ~(u64)(SEMA_PAGE_SIZE - 1)) -
+ (group_pa - s);
+ ncand++;
+ }
}
}
return ncand;
@@ -6175,52 +6997,85 @@ static bool svc_rescue_payload_ok(u32 slot, u32 want, const u8 *buf, u32 len)
{
u64 dummy[SVC_MBUF_MAX];
- return svc_rescue_match_buf(buf, len, 0, slot, want, len,
- dummy, 0) > 0;
+ return svc_rescue_match_buf(buf, len, 0, slot, want, 0, len,
+ dummy, NULL, 0, NULL) > 0;
}
/*
- * Pin a candidate page as the translation of the entry VA's page and verify
- * by reading the whole push back and parsing it as a method stream. Requires
- * the push to end exactly on a method boundary - the stream replaces a weak
- * prefix check that accepted page-table pages (whose first word can decode as
- * a header). Returns 0 and leaves the pin installed on success.
+ * Verify a candidate page as the translation of one VA page of the entry's
+ * push. The whole push is assembled by *direct physical reads*: the
+ * candidate's VA page reads from cand_pa, every other page through the
+ * existing trusted translations. Nothing is written to the pin table until
+ * the assembled push parses as a complete method stream and carries the
+ * expected release, so a rejected candidate cannot poison the very pages it
+ * was tried for (observed: failed rescue attempts overwrote previously
+ * validated pins with window guesses, turning readable pool pages into
+ * zeros and cascading stalls). On success the candidate is pinned. Returns 0
+ * on success, -ENOENT on validation failure.
*/
static int svc_rescue_try(struct nv_vgpu_device *vgpu, struct svc_channel *ch,
- u64 entry_va, u32 entry_len, u32 pin_len, u64 cand_pa,
- u32 cand_i, u32 ncand)
+ u64 entry_va, u32 entry_len, u64 cand_pa,
+ u64 cand_va_page, u32 cand_i, u32 ncand)
{
u8 *full;
- u32 j;
- int rc = -ENOENT;
-
- for (j = 0; j < pin_len; j += SEMA_PAGE_SIZE) {
- svc_pin_store(entry_va + j, cand_pa + j, SVC_APERTURE_SYS);
- svc_lat_store(entry_va + j, cand_pa + j, SVC_APERTURE_SYS);
- }
+ u32 off = 0;
+ int rc = 0;
full = kmalloc(entry_len, GFP_KERNEL);
if (!full)
return -ENOMEM;
- rc = svc_read_va_trusted(vgpu, entry_va, full, entry_len);
+ while (off < entry_len) {
+ u64 va = entry_va + off;
+ size_t chunk = min_t(size_t, entry_len - off,
+ SEMA_PAGE_SIZE -
+ (va & (SEMA_PAGE_SIZE - 1)));
+ u64 pa;
+ int ap = SVC_APERTURE_SYS;
+
+ if ((va & ~(u64)(SEMA_PAGE_SIZE - 1)) == cand_va_page) {
+ pa = (cand_pa & ~(u64)(SEMA_PAGE_SIZE - 1)) |
+ (va & (SEMA_PAGE_SIZE - 1));
+ } else if (svc_lookup_trusted(vgpu, va, &pa, &ap)) {
+ rc = -ENOENT;
+ break;
+ }
+ if (svc_read_ap(vgpu, pa,
+ ap == SVC_APERTURE_FB ? SVC_APERTURE_FB :
+ SVC_APERTURE_SYS,
+ full + off, chunk)) {
+ rc = -EIO;
+ break;
+ }
+ off += chunk;
+ }
if (!rc && svc_stream_full_ok(full, entry_len) &&
svc_rescue_payload_ok(4 * ch->ce_index,
ch->expected_sem_payload, full,
entry_len)) {
- for (j = 0; j < pin_len; j += SEMA_PAGE_SIZE)
- svc_mark_hot(cand_pa + j);
- pr_info("sema-svc: content rescue ch=0x%x rank=%u slot=%u entry_va=0x%llx len=%u want=%u -> pa=0x%llx (cand %u/%u)\n",
+ svc_pin_store(cand_va_page, cand_pa, SVC_APERTURE_SYS);
+ svc_lat_store(cand_va_page, cand_pa, SVC_APERTURE_SYS);
+ svc_mark_hot(cand_pa);
+ pr_info("sema-svc: content rescue ch=0x%x rank=%u slot=%u entry_va=0x%llx len=%u want=%u -> pa=0x%llx va_page=0x%llx (cand %u/%u)\n",
ch->h_object, ch->ce_index, 4 * ch->ce_index, entry_va,
- entry_len, ch->expected_sem_payload, cand_pa, cand_i + 1, ncand);
+ entry_len, ch->expected_sem_payload, cand_pa,
+ cand_va_page, cand_i + 1, ncand);
rc = 0;
} else if (gsp_sema_dbg && svc.rescue_rej_log < 32) {
- u32 w0 = 0;
+ u32 w0 = 0, craw = 0;
+ u64 ipa = 0;
+ int iap = SVC_APERTURE_SYS;
+ bool have_ipin;
svc.rescue_rej_log++;
if (!rc && entry_len >= 4)
memcpy(&w0, full, 4);
- pr_info("sema-svc: rescue reject cand=0x%llx va=0x%llx read_rc=%d w0=0x%08x\n",
- cand_pa, entry_va, rc, w0);
+ svc_read(vgpu, (cand_pa & ~(u64)(SEMA_PAGE_SIZE - 1)) |
+ (entry_va & (SEMA_PAGE_SIZE - 1)),
+ &craw, sizeof(craw));
+ have_ipin = !svc_ipin_lookup(entry_va, &ipa, &iap);
+ pr_info("sema-svc: rescue reject cand=0x%llx va=0x%llx va_page=0x%llx read_rc=%d w0=0x%08x cand_first=0x%08x ipin=%d ipa=0x%llx want=%u\n",
+ cand_pa, entry_va, cand_va_page, rc, w0, craw,
+ have_ipin ? 1 : 0, ipa, ch->expected_sem_payload);
rc = -ENOENT;
}
kfree(full);
@@ -6234,8 +7089,9 @@ static int svc_rescue_try(struct nv_vgpu_device *vgpu, struct svc_channel *ch,
*/
static int svc_rescue_scan_range(struct nv_vgpu_device *vgpu,
struct svc_channel *ch, u64 entry_va,
- u32 entry_len, u32 pin_len, u32 slot,
- u64 hi, u64 lo, u8 *buf, u64 *cand)
+ u32 entry_len, u32 slot,
+ u64 hi, u64 lo, u8 *buf, u64 *cand,
+ u64 *cand_va)
{
u64 gpa = hi;
unsigned long t0 = jiffies;
@@ -6250,12 +7106,14 @@ static int svc_rescue_scan_range(struct nv_vgpu_device *vgpu,
ncand = svc_rescue_match_buf(buf, SVC_SCAN_CHUNK, gpa,
slot,
ch->expected_sem_payload,
- entry_len, cand, 0);
+ entry_va, entry_len,
+ cand, cand_va, 0,
+ &ch->content_max);
total += ncand;
for (n = 0; n < ncand; n++) {
if (!svc_rescue_try(vgpu, ch, entry_va,
- entry_len, pin_len,
- cand[n], n, ncand))
+ entry_len, cand[n],
+ cand_va[n], n, ncand))
return 0;
}
}
@@ -6285,9 +7143,8 @@ static int svc_rescue_push_page(struct nv_vgpu_device *vgpu,
u8 *buf;
u64 gpa;
u32 slot = 4 * ch->ce_index;
- u64 cand[SVC_MBUF_MAX];
+ u64 cand[SVC_MBUF_MAX], cand_va[SVC_MBUF_MAX];
u32 ncand = 0, i;
- u32 pin_len;
/*
* The matcher needs only the channel rank (slot) and the release-push
@@ -6305,16 +7162,6 @@ static int svc_rescue_push_page(struct nv_vgpu_device *vgpu,
return -EINVAL;
/*
- * Slot-aligned entries pin the whole 1 MB slot at once (all of its
- * pushes share the mapping); an entry in the middle of a slot only
- * needs the pages its own push occupies - the page-walking candidates
- * in svc_rescue_match_buf cover the push starting on a page before the
- * release.
- */
- pin_len = (entry_va & 0xfffff) ?
- ALIGN(entry_len + SEMA_PAGE_SIZE, SEMA_PAGE_SIZE) : 0x100000;
-
- /*
* Fast path: once a channel's push-buffer window base is verified
* (svc_pb_pair pinned it from the first push), a later entry's page
* sits at the same window-relative offset in guest RAM. Validate the
@@ -6329,8 +7176,9 @@ static int svc_rescue_push_page(struct nv_vgpu_device *vgpu,
entry_va < svc_pb[i].va)
continue;
cand_pa = svc_pb[i].pa_page + (entry_va - svc_pb[i].va);
- if (!svc_rescue_try(vgpu, ch, entry_va, entry_len, pin_len,
- cand_pa, i, svc_pb_n))
+ if (!svc_rescue_try(vgpu, ch, entry_va, entry_len, cand_pa,
+ entry_va & ~(u64)(SEMA_PAGE_SIZE - 1),
+ i, svc_pb_n))
return 0;
}
@@ -6359,11 +7207,13 @@ static int svc_rescue_push_page(struct nv_vgpu_device *vgpu,
continue;
ncand = svc_rescue_match_buf(buf, SEMA_PAGE_SIZE, gpa,
slot, ch->expected_sem_payload,
- entry_len, cand, 0);
+ entry_va, entry_len, cand,
+ cand_va, 0,
+ &ch->content_max);
for (n = 0; n < ncand; n++) {
if (!svc_rescue_try(vgpu, ch, entry_va,
- entry_len, pin_len,
- cand[n], n, ncand)) {
+ entry_len, cand[n],
+ cand_va[n], n, ncand)) {
kfree(buf);
return 0;
}
@@ -6410,7 +7260,7 @@ static int svc_rescue_push_page(struct nv_vgpu_device *vgpu,
hi = SEMA_PA_MAX - SVC_SCAN_CHUNK;
if (lo <= hi &&
!svc_rescue_scan_range(vgpu, ch, entry_va, entry_len,
- pin_len, slot, hi, lo, buf, cand)) {
+ slot, hi, lo, buf, cand, cand_va)) {
kfree(buf);
return 0;
}
@@ -6431,11 +7281,43 @@ static int svc_rescue_push_page(struct nv_vgpu_device *vgpu,
* 4.4+ GB to reach them, a descending one finds them in the first
* few hundred MB.
*/
- i = svc_rescue_scan_range(vgpu, ch, entry_va, entry_len, pin_len, slot,
- SEMA_PA_MAX - SVC_SCAN_CHUNK, 0, buf, cand);
+ i = svc_rescue_scan_range(vgpu, ch, entry_va, entry_len, slot,
+ SEMA_PA_MAX - SVC_SCAN_CHUNK, 0, buf, cand,
+ cand_va);
kfree(buf);
if (!i)
return 0;
+
+ /*
+ * Content watermark re-sync. A fake release advances the
+ * expectation even when its write is only queued, so a burst of
+ * unreadable entries can push `expected` past every push the guest
+ * has actually written (observed: expected=67 while RAM's newest
+ * release for the slot was 53; the exact-payload matcher then
+ * rejects the real pages and the scan can never recover). The
+ * highest payload the scans saw for this slot is the real
+ * watermark: when it sits behind the expectation, re-sync down to
+ * watermark+1 so the next attempt demands the payload the guest's
+ * next push actually carries. Forward-only, never below the
+ * watermark.
+ *
+ * The drift a fake burst can create is small (the number of
+ * entries consumed while content was unreadable); a watermark
+ * far behind the expectation is old RAM content, not the push
+ * watermark, and must not drag the expectation back to it.
+ */
+ if (ch->content_max && ch->content_max + 1 < ch->expected_sem_payload &&
+ ch->expected_sem_payload - ch->content_max <= 64) {
+ u32 want = ch->content_max + 1;
+
+ if (svc.resc_resync < 64)
+ pr_info("sema-resync: h=0x%x rank=%u expected=%u -> %u (content watermark=%u)\n",
+ ch->h_object, ch->ce_index,
+ ch->expected_sem_payload, want,
+ ch->content_max);
+ svc.resc_resync++;
+ ch->expected_sem_payload = want;
+ }
return -ENOENT;
}
@@ -7286,7 +8168,9 @@ static void svc_fwd_hold_expire_check(struct nv_vgpu_device *vgpu,
struct svc_channel *ch)
{
struct nv_vgpu_device *hv;
- struct svc_channel shim;
+ static struct svc_channel shim; /* single-threaded service path; BSS to
+ * keep the 16 KB kernel stack safe
+ * after the channel struct grew */
bool done;
if (!READ_ONCE(svc_fwd_hold.active) ||
@@ -7462,7 +8346,9 @@ static void svc_fwd_hold_gone_finish(struct nv_vgpu_device *vgpu)
{
struct nv_vgpu_device *hv = svc_fwd_hold.vgpu ?
svc_fwd_hold.vgpu : vgpu;
- struct svc_channel shim;
+ static struct svc_channel shim; /* single-threaded service path; BSS to
+ * keep the 16 KB kernel stack safe
+ * after the channel struct grew */
bool done = smp_load_acquire(&svc_fwd_req.done);
bool aborted = READ_ONCE(svc_fwd_hold.aborted);
@@ -7511,7 +8397,9 @@ static void svc_fwd_hold_stop_flush(struct nv_vgpu_device *vgpu)
{
struct nv_vgpu_device *hv = svc_fwd_hold.vgpu ?
svc_fwd_hold.vgpu : vgpu;
- struct svc_channel shim;
+ static struct svc_channel shim; /* single-threaded service path; BSS to
+ * keep the 16 KB kernel stack safe
+ * after the channel struct grew */
bool done;
if (!svc_fwd_hold.active)
@@ -8207,6 +9095,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
u8 entry[SVC_GPFIFO_ENTRY_SIZE];
u64 va, idx = ch->last_get;
u32 len, e0, e1, g;
+ u64 entry_ident = 0;
+ bool entry_replay = false;
/*
* P6b wave 1.5: the held entry is resolved before
@@ -8275,12 +9165,73 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
/* GP_ENTRY1.LENGTH[30:10], in dwords */
len = ((e1 >> 10) & 0x1fffff) << 2;
- if (gsp_sema_dbg && svc.gpfifo_entry_log < 32) {
- pr_info("sema-chan: entry[%llu] raw=%08x%08x va=0x%llx len=%u\n",
- idx, e1, e0, va, len);
+ /*
+ * Entry identity for once-only op execution: a pin
+ * rewind replays consumed entries, and a CE copy read
+ * from a client staging buffer that has since been
+ * refilled must not run again (it would copy the next
+ * chunk's data over the finished one). Releases still
+ * re-apply below; only the ops are gated.
+ */
+ {
+ u64 ident = (u64)e0 | ((u64)e1 << 32);
+ bool replay_entry = ident &&
+ ch->slot_ops_ident[idx] == ident;
+
+ entry_ident = ident;
+ entry_replay = replay_entry;
+ }
+
+ if (gsp_sema_dbg && svc.gpfifo_entry_log < 160) {
+ pr_info("sema-chan: entry[%llu] h=0x%x eng=%u raw=%08x%08x va=0x%llx len=%u\n",
+ idx, ch->h_object, ch->engine_type,
+ e1, e0, va, len);
svc.gpfifo_entry_log++;
}
+ /*
+ * Cheap wait for an entry already inside its
+ * expensive-wait window: read + rescue + hold were
+ * already paid for this VA, and a pin-batch rewind
+ * replays entries every few milliseconds while new
+ * mappings stream in. Repaying the work per replay
+ * is a livelock (observed: 223 replays of one
+ * 12.8 KB entry in 9 s, starving the other
+ * channels and flooding the log). Wait cheaply
+ * until a fresh injected/validated translation
+ * exists for the page, the put advances (a new
+ * push at the same VA), or the window expires.
+ */
+ if (ch->engine_type && ch->unread_va == va &&
+ ch->unread_idx == idx &&
+ time_before(jiffies, ch->unread_until)) {
+ u64 tpa = 0, vpa = 0;
+ int tap = 0, vap = 0;
+ bool have_ipin;
+ bool have_val;
+ bool fresh;
+
+ /*
+ * Both lookups return 0 on success. Treat a
+ * validated pin, or an injected pin that
+ * differs from the one the window opened
+ * with, as fresh evidence and fall through
+ * to the normal read/rescue path; otherwise
+ * wait cheaply. The same (possibly wrong)
+ * injection re-appearing must not re-run
+ * the rescue on every pin-batch rewind.
+ */
+ have_ipin = svc_ipin_lookup(va, &tpa, &tap) == 0;
+ have_val = svc_pin_lookup_valid(va, &vpa, &vap) == 0;
+ fresh = have_val ||
+ (have_ipin && tpa != ch->unread_ipin_pa);
+ if (!fresh) {
+ svc.gpfifo_held++;
+ svc.gpfifo_nomap++;
+ break;
+ }
+ }
+
if (va && len && len <= SVC_PUSH_MAX) {
u8 *pb = kmalloc(len, GFP_KERNEL);
@@ -8348,7 +9299,45 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
}
}
+ /*
+ * Content rescue before any fake-release
+ * fallback (CE channels): the push bytes
+ * are in guest RAM even when the VA->PA
+ * translation is missing. A candidate is
+ * accepted only when the whole entry
+ * parses as a method stream and carries
+ * this channel's exact next release (slot
+ * + payload); the page is then pinned and
+ * the push read and executed like a
+ * trusted push. Without this the pump
+ * acknowledged the release while the
+ * copy that stages the data had never
+ * run - the guest read unstaged memory
+ * (observed: cuMemcpyDtoH returning
+ * zeros, noclaim8). The probe is
+ * synchronous and bounded (hot set ->
+ * near scan -> rate-limited fallback);
+ * on failure the entry follows the old
+ * fake/hold path unchanged, so it can
+ * never livelock the channel.
+ */
+ if (unreadable && ch->engine_type &&
+ ch->expected_sem_payload) {
+ if (!svc_rescue_push_page(vgpu, ch,
+ va, len) &&
+ !svc_read_va_push(vgpu, ch, va,
+ pb, len) &&
+ svc_method_stream_plausible(pb,
+ len)) {
+ ch->rescue_last = 0;
+ ch->push_hold_until = 0;
+ svc.gpfifo_rescued++;
+ unreadable = false;
+ }
+ }
+
if (!unreadable) {
+ ch->push_hold_until = 0;
/*
* Fills/copies execute only for CE
* channels: the push was written for the
@@ -8365,7 +9354,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
svc.op_tr_pending = 0;
svc_decode_stream(vgpu, pb,
len, 0,
- allow_ops || injected,
+ (allow_ops || injected) &&
+ !entry_replay,
ch,
idx,
va,
@@ -8410,6 +9400,15 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
break;
}
ch->op_retry = 0;
+ /*
+ * Ops (if any) have now run for
+ * this entry identity; remember it
+ * so a rewind replay cannot run
+ * them a second time.
+ */
+ if (entry_ident)
+ ch->slot_ops_ident[idx] =
+ entry_ident;
rel_seen = (svc_rel_count != rel_n_before);
{
bool fwd_guest_queued;
@@ -8517,17 +9516,74 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
ch->expected_sem_payload) {
u32 rel_n2 = svc_rel_count;
- if (!svc_rescue_push_page(vgpu, ch, va, len) &&
- !svc_read_va_push(vgpu, ch, va, pb, len))
- svc_decode_stream(vgpu, pb,
- len, 0, false,
- ch, idx, va, true);
- rel_seen =
- (svc_rel_count != rel_n2);
+ /*
+ * Open the cheap-wait window
+ * for this entry (see the
+ * loop-top guard): the rescue
+ * below is the expensive work
+ * a rewind must not repay.
+ */
+ ch->unread_va = va;
+ ch->unread_idx = (u32)idx;
+ if (!ch->unread_until) {
+ u64 ipa = 0;
+ int iap = 0;
+
+ ch->unread_until =
+ jiffies +
+ msecs_to_jiffies(gsp_sema_push_hold_ms ?
+ gsp_sema_push_hold_ms :
+ gsp_sema_ce_grace_ms);
+ ch->unread_ipin_pa = 0;
+ if (svc_ipin_lookup(va, &ipa,
+ &iap) == 0)
+ ch->unread_ipin_pa = ipa;
+ }
+ {
+ u32 want_before =
+ ch->expected_sem_payload;
+
+ if (!svc_rescue_push_page(vgpu, ch, va, len) &&
+ !svc_read_va_push(vgpu, ch, va, pb, len))
+ svc_decode_stream(vgpu, pb,
+ len, 0, false,
+ ch, idx, va, true);
+ rel_seen =
+ (svc_rel_count != rel_n2);
+ /*
+ * The rescue re-synced the
+ * expectation down to the RAM
+ * content watermark (the fakes
+ * had pushed it past every push
+ * the guest actually wrote):
+ * retry once, now demanding the
+ * payload the content shows.
+ * Without this the pump chases
+ * the watermark one release per
+ * full scan and never reaches
+ * the entry the client waits
+ * on.
+ */
+ if (!rel_seen &&
+ ch->expected_sem_payload !=
+ want_before) {
+ ch->rescue_last = 0;
+ if (!svc_rescue_push_page(vgpu, ch, va, len) &&
+ !svc_read_va_push(vgpu, ch, va, pb, len))
+ svc_decode_stream(vgpu, pb,
+ len, 0, false,
+ ch, idx, va, true);
+ rel_seen =
+ (svc_rel_count != rel_n2);
+ }
+ }
}
- if (rel_seen)
+ if (rel_seen) {
ch->last_rel_va =
svc_rel_va_seen;
+ ch->unread_va = 0;
+ ch->unread_until = 0;
+ }
ch->submissions++;
svc.gpfifo_entries_done++;
if (ch->engine_type && ch->expected_sem_payload) {
@@ -8549,6 +9605,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
ch->entry_want) {
ch->defer_since = 0;
ch->rescue_fail = 0;
+ ch->unread_va = 0;
+ ch->unread_until = 0;
svc_replay_note(ch, idx, put);
ch->last_get = (idx + 1) %
ch->gpfifo_entries;
@@ -8594,6 +9652,65 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
*/
svc_log_stalled_va(va);
/*
+ * Open the cheap-wait window
+ * for this entry (see the
+ * loop-top guard) so rewind
+ * replays do not repay the
+ * rescue on every pin batch.
+ */
+ ch->unread_va = va;
+ ch->unread_idx = (u32)idx;
+ if (!ch->unread_until) {
+ u64 ipa = 0;
+ int iap = 0;
+
+ ch->unread_until =
+ jiffies +
+ msecs_to_jiffies(gsp_sema_push_hold_ms ?
+ gsp_sema_push_hold_ms :
+ gsp_sema_ce_grace_ms);
+ ch->unread_ipin_pa = 0;
+ if (svc_ipin_lookup(va, &ipa,
+ &iap) == 0)
+ ch->unread_ipin_pa = ipa;
+ }
+ /*
+ * Bounded work-before-release
+ * hold (CE): the deferred retry
+ * probe shows the page's
+ * translation/pins land
+ * 0.5-2 s after the entry is
+ * first seen (sema-retry:
+ * "resolvable after ..."). Hold
+ * the entry for a fixed absolute
+ * grace so a late translation
+ * executes the work before the
+ * release reaches the client.
+ * The deadline is never reset by
+ * the replay/rescue paths and is
+ * cleared when the entry is
+ * consumed, so it cannot
+ * livelock (the noclaim9
+ * lesson); on expiry the entry
+ * falls through to the fake
+ * fallback exactly as before.
+ */
+ if (gsp_sema_push_hold_ms &&
+ ch->engine_type &&
+ ch->expected_sem_payload) {
+ if (!ch->push_hold_until)
+ ch->push_hold_until =
+ jiffies +
+ msecs_to_jiffies(gsp_sema_push_hold_ms);
+ if (time_before(jiffies,
+ ch->push_hold_until)) {
+ svc.gpfifo_held++;
+ svc.gpfifo_nomap++;
+ kfree(pb);
+ break;
+ }
+ }
+ /*
* The push's work (the data copies
* the client reads back) must run
* before the release the client
@@ -8631,7 +9748,18 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
}
}
if (svc_fake_release(vgpu, ch)) {
+ /*
+ * Watch (observe only):
+ * did this push's page
+ * ever become readable?
+ * See svc_retry_poll().
+ */
+ svc_retry_enqueue(va, len,
+ ch->expected_sem_payload);
+ ch->push_hold_until = 0;
ch->defer_since = 0;
+ ch->unread_va = 0;
+ ch->unread_until = 0;
svc_replay_note(ch, idx, put);
ch->last_get = (idx + 1) %
ch->gpfifo_entries;
@@ -8647,6 +9775,7 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
ch, va, len);
if (!rescue_rc) {
ch->rescue_fail = 0;
+ ch->rescue_last = 0;
/* Rescue succeeded, retry the read. */
if (!svc_read_va_push(vgpu, ch, va, pb,
len)) {
@@ -8723,6 +9852,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
*/
ch->rescue_fail = 0;
ch->defer_since = 0;
+ ch->unread_va = 0;
+ ch->unread_until = 0;
svc_replay_note(ch, idx, put);
ch->last_get =
(idx + 1) %
@@ -8825,6 +9956,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
svc.gpfifo_starved++;
ch->rescue_fail = 0;
ch->defer_since = 0;
+ ch->unread_va = 0;
+ ch->unread_until = 0;
if (gsp_sema_fake_rel ||
gsp_sema_rel_recover)
svc_replay_note(ch, idx,
@@ -8869,6 +10002,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
ch->last_get = (idx + 1) % ch->gpfifo_entries;
ch->rescue_fail = 0;
ch->defer_since = 0;
+ ch->unread_va = 0;
+ ch->unread_until = 0;
/* Publish progress the way the GPU would. */
g = ch->last_get;
@@ -8925,9 +10060,25 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
svc.chan[k].last_rel_va = ch->last_rel_va;
svc.chan[k].rescue_last = ch->rescue_last;
svc.chan[k].rescue_fail = ch->rescue_fail;
+ svc.chan[k].push_hold_until = ch->push_hold_until;
svc.chan[k].defer_since = ch->defer_since;
svc.chan[k].put_prev = ch->put_prev;
/*
+ * Persist the cheap-wait window and the content
+ * watermark: the pump works on a snapshot, so
+ * without the write-back the rewind guard would
+ * never see a window and each pass would repay
+ * the read+rescue for an entry pin-batch replays
+ * are re-opening.
+ */
+ svc.chan[k].unread_va = ch->unread_va;
+ svc.chan[k].unread_idx = ch->unread_idx;
+ svc.chan[k].unread_until = ch->unread_until;
+ svc.chan[k].unread_ipin_pa = ch->unread_ipin_pa;
+ svc.chan[k].pushpg_va = ch->pushpg_va;
+ svc.chan[k].pushpg_n = ch->pushpg_n;
+ svc.chan[k].content_max = ch->content_max;
+ /*
* P6c C3: the delivered-fault notifier latch
* lives on the snapshot copy; persist it so the
* next pass's svc_notify_complete() keeps
@@ -8942,6 +10093,14 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
ch->slot_done_epoch,
SVC_GPFIFO_MAX_BITS);
/*
+ * Persist the once-only op-execution identities:
+ * the pump works on a snapshot, so without this
+ * a rewind replay would re-run consumed copies.
+ */
+ memcpy(svc.chan[k].slot_ops_ident,
+ ch->slot_ops_ident,
+ sizeof(ch->slot_ops_ident));
+ /*
* Persist the counter slot autoreg resolved on the
* snapshot: without this the next pass re-derives
* it (with a different live GR count during
@@ -9590,6 +10749,12 @@ static int svc_thread(void *arg)
*/
svc_apply_pushes(vgpu);
+ /*
+ * Deferred readability probes for faked CE pushes (observation
+ * only; see svc_retry_poll).
+ */
+ svc_retry_poll(vgpu);
+
msleep(10);
}
@@ -9706,6 +10871,8 @@ void gsp_sema_service_stop(void)
}
/* Thread is gone: release the bridge's pending jobs and debugfs files. */
gsp_cuda_bridge_fini();
+ /* A fresh session must not inherit shadowed bytes from the last one. */
+ svc_vas_free_all();
}
/* ------------------------------------------------------------------ */
@@ -9713,10 +10880,11 @@ void gsp_sema_service_stop(void)
/* ------------------------------------------------------------------ */
/*
- * Page-wise GPU-VA read for the bridge: resolve each page through the same
- * trusted sources every operand read uses (pins, learned cache, discovered
- * map, published page directories, targeted rm_mem resolve) and read through
- * the aperture the VA maps to (vidmem shadow vs guest RAM via vfio_dma_rw).
+ * Page-wise GPU-VA read for the bridge: a page the VA shadow already holds
+ * is served from it, otherwise resolve through the same trusted sources
+ * every operand read uses (pins, learned cache, discovered map, published
+ * page directories, targeted rm_mem resolve) and read through the aperture
+ * the VA maps to (vidmem shadow vs guest RAM via vfio_dma_rw).
*/
int gsp_sema_service_va_read(struct nv_vgpu_device *vgpu, u64 va, void *buf,
size_t len)
@@ -9730,10 +10898,15 @@ int gsp_sema_service_va_read(struct nv_vgpu_device *vgpu, u64 va, void *buf,
u64 addr = 0;
int ap = SVC_APERTURE_SYS;
- if (svc_lookup(vgpu, va, &addr, &ap))
- return -ENOENT;
- if (svc_read_ap(vgpu, addr, svc_ap_of(ap), p, chunk))
- return -EIO;
+ if (svc_vas_has(va)) {
+ if (svc_vas_read(va, p, chunk))
+ return -EIO;
+ } else {
+ if (svc_io_resolve(vgpu, va, &addr, &ap))
+ return -ENOENT;
+ if (svc_read_ap(vgpu, addr, svc_ap_of(ap), p, chunk))
+ return -EIO;
+ }
p += chunk;
va += chunk;
len -= chunk;
@@ -9741,6 +10914,14 @@ int gsp_sema_service_va_read(struct nv_vgpu_device *vgpu, u64 va, void *buf,
return 0;
}
+/*
+ * Page-wise GPU-VA write for the bridge. The bytes always land in the VA
+ * shadow (so shadow-served reads and a later flush are consistent), and
+ * additionally write through the translation when one exists: the guest's
+ * read-back of a delivered buffer may be a CPU read of a guest page the
+ * translation names, not an emulated copy, so shadow-only delivery is not
+ * enough.
+ */
int gsp_sema_service_va_write(struct nv_vgpu_device *vgpu, u64 va,
const void *buf, size_t len)
{
@@ -9753,10 +10934,22 @@ int gsp_sema_service_va_write(struct nv_vgpu_device *vgpu, u64 va,
u64 addr = 0;
int ap = SVC_APERTURE_SYS;
- if (svc_lookup(vgpu, va, &addr, &ap))
- return -ENOENT;
- if (svc_write_ap(vgpu, addr, svc_ap_of(ap), (void *)p, chunk))
- return -EIO;
+ if (svc_vas_write(va, p, chunk))
+ return -ENOMEM;
+ if (!svc_io_resolve(vgpu, va, &addr, &ap)) {
+ if (svc_write_ap(vgpu, addr, svc_ap_of(ap),
+ (void *)p, chunk))
+ return -EIO;
+ if (!svc.vas_flushed) {
+ svc.vas_flushed = 1;
+ pr_info("sema-vas: bridge write flushed through translation va=0x%llx pa=0x%llx ap=%s\n",
+ va, addr,
+ svc_ap_of(ap) == SVC_APERTURE_FB ?
+ "vidmem" : "ram");
+ }
+ } else {
+ svc.vas_shadowed++;
+ }
p += chunk;
va += chunk;
len -= chunk;
@@ -9764,14 +10957,46 @@ int gsp_sema_service_va_write(struct nv_vgpu_device *vgpu, u64 va,
return 0;
}
+/*
+ * Translation-only view for the bridge: is there a PA for this VA right now,
+ * ignoring the shadow? Used by the completion path to wait (bounded) for the
+ * guest to publish the mapping instead of accepting a shadow-only delivery
+ * the guest cannot read.
+ */
+int gsp_sema_service_va_translate(struct nv_vgpu_device *vgpu, u64 va,
+ u64 *pa, u32 *ap)
+{
+ u64 addr = 0;
+ int iap = SVC_APERTURE_SYS;
+
+ if (svc_io_resolve(vgpu, va, &addr, &iap))
+ return -ENOENT;
+ if (pa)
+ *pa = addr;
+ if (ap)
+ *ap = svc_ap_of(iap);
+ return 0;
+}
+
int gsp_sema_service_va_resolve(struct nv_vgpu_device *vgpu, u64 va, u64 *pa,
u32 *ap)
{
u64 addr = 0;
int iap = SVC_APERTURE_SYS;
- if (svc_lookup(vgpu, va, &addr, &iap))
+ if (svc_io_resolve(vgpu, va, &addr, &iap)) {
+ /* A shadowed page is resolvable for the data path even though
+ * it has no PA: report success with the shadow bit set so the
+ * bridge's diagnostics do not read as a data-path failure. */
+ if (svc_vas_has(va)) {
+ if (pa)
+ *pa = 0;
+ if (ap)
+ *ap = 0xff; /* SVC_APERTURE_VASHDW */
+ return 0;
+ }
return -ENOENT;
+ }
if (pa)
*pa = addr;
if (ap)
@@ -9779,6 +11004,56 @@ int gsp_sema_service_va_resolve(struct nv_vgpu_device *vgpu, u64 va, u64 *pa,
return 0;
}
+/*
+ * One-shot walk explainer: dump every level of every published page
+ * directory (and the self-discovered one) for a VA whose translation keeps
+ * failing, so a run records why the walk misses even when the VA-shadow path
+ * rescues the data. Purely diagnostic; no state changes beyond lat cache.
+ */
+void gsp_sema_service_va_explain(struct nv_vgpu_device *vgpu, u64 va)
+{
+ u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1);
+ int i;
+
+ pr_info("sema-explain: va=0x%llx page=0x%llx shadow=%d discovered=%d base=0x%llx entries=%u\n",
+ va, page, svc_vas_has(va), svc_pd_found, svc_pd_base,
+ svc_pd_entries);
+ for (i = 0; i < SVC_MAX_VASPACES; i++) {
+ u64 a = 0;
+ int ap = 0;
+ int rc;
+
+ if (!svc.vas[i].valid)
+ continue;
+ pr_info("sema-explain: vas[%d] h=0x%x pdb=0x%llx entries=%u ap=%d\n",
+ i, svc.vas[i].h_vaspace, svc.vas[i].pdb,
+ svc.vas[i].entries, svc.vas[i].ap);
+ rc = svc_gmmu_walk(vgpu, svc.vas[i].pdb, svc.vas[i].entries,
+ page, svc.vas[i].ap, &a, &ap, true);
+ pr_info("sema-explain: vas[%d] walk rc=%d addr=0x%llx ap=%d\n",
+ i, rc, a, ap);
+ }
+ if (svc_pd_found) {
+ u64 a = 0;
+ int ap = 0;
+ int rc = svc_gmmu_walk(vgpu, svc_pd_base, svc_pd_entries, page,
+ svc_pd_ap, &a, &ap, true);
+
+ pr_info("sema-explain: discovered walk rc=%d addr=0x%llx ap=%d\n",
+ rc, a, ap);
+ }
+}
+
+/*
+ * Diagnostic window: for `ms` milliseconds after a bridge completion, log
+ * the guest's decoded operations (pushes, LAUNCH_DMAs, ops) so the read-back
+ * encoding of a delivered buffer is on record. Purely observational.
+ */
+void gsp_sema_service_trace_window(u32 ms)
+{
+ svc.op_trace_until = jiffies + msecs_to_jiffies(ms);
+}
+
int gsp_sema_service_sem_read(struct nv_vgpu_device *vgpu, u64 va, u32 *val)
{
return svc_sem_read(vgpu, va, val);
diff --git a/kernel/gsp_sema_service.h b/kernel/gsp_sema_service.h
index 530663e..22e9edc 100644
--- a/kernel/gsp_sema_service.h
+++ b/kernel/gsp_sema_service.h
@@ -210,6 +210,13 @@ extern u64 gsp_sema_copy_dst2;
extern unsigned gsp_sema_ce_grace_ms;
/*
+ * Bounded work-before-release hold (ms) for an unreadable CE entry whose
+ * translation/pins typically land 0.5-2 s late (observed via the deferred
+ * retry probe). Absolute deadline per entry, never reset; 0 disables.
+ */
+extern unsigned gsp_sema_push_hold_ms;
+
+/*
* Grace window (ms) an untranslatable GR GPFIFO entry is left pending before
* the pump consumes it as stale. The GR push pages are pinned from guest
* truth a moment after the channels appear; the window lets those pins land
diff --git a/kernel/include/nvidia-vgpu-vfio.h b/kernel/include/nvidia-vgpu-vfio.h
index 1c2f67b..ccd7262 100644
--- a/kernel/include/nvidia-vgpu-vfio.h
+++ b/kernel/include/nvidia-vgpu-vfio.h
@@ -25,7 +25,18 @@
#define NV_VGPU_MAX_IRQS NVIDIA_VGPU_MAX_IRQS
#define NV_VGPU_BAR0_SIZE (16 * 1024 * 1024)
-#define NV_VGPU_BAR1_SIZE (256 * 1024 * 1024)
+/*
+ * BAR1 (VRAM aperture) and the emulated framebuffer size. These two move
+ * together: the guest derives its BAR1 size from fbLength
+ * (kbusInitBarsSize_VF: pciBarSizes[1] = pGpu->fbLength), and its vidmem
+ * address space is [0, fbLength). The shadow backing bars[1].vaddr is
+ * NV_VGPU_BAR1_SIZE, so any guest vidmem PA the emulation translates must
+ * stay inside it. 512 MiB gives UVM enough headroom after its own
+ * reservations to satisfy the acceptance program's 3x4 MiB vadd buffers
+ * (at 256 MiB the guest reported only ~5 MiB free).
+ */
+#define NV_VGPU_BAR1_SIZE (512 * 1024 * 1024)
+#define NV_VGPU_EMUL_FB_SIZE (512ULL * 1024 * 1024)
#define NV_VGPU_BAR2_SIZE (32 * 1024 * 1024)
#define NV_VGPU_BAR3_SIZE (32 * 1024 * 1024)
@@ -56,6 +67,13 @@ struct nv_vgpu_bar {
void *vaddr;
bool mapped;
bool trapped;
+ /*
+ * True when vaddr is a host-RAM shadow (vzalloc) owned by this module
+ * instead of an ioremap'd physical BAR. Guest mappings must then remap
+ * the shadow's own pages; a real BAR must never be exposed to the
+ * guest outside the supervised physical claim.
+ */
+ bool emulated;
struct mutex lock;
};
diff --git a/kernel/main.c b/kernel/main.c
index 7001b3d..ca973b0 100644
--- a/kernel/main.c
+++ b/kernel/main.c
@@ -1285,15 +1285,34 @@ static int nv_vgpu_vfio_dev_mmap(struct vfio_device *vdev, struct vm_area_struct
return 0;
}
- /* BAR2 region - ioremap'd host BAR; map backing physical address */
+ /* BAR2 region - emulated shadow (emulation-only) or real BAR (claim) */
region_base = VFIO_PCI_INDEX_TO_OFFSET(VFIO_PCI_BAR2_REGION_INDEX);
region_size = vgpu->bars[2].size;
if (offset >= region_base && offset < region_base + region_size) {
+ unsigned long bar2_off = offset - region_base;
+ unsigned long pos;
+
if (!vgpu->bars[2].vaddr || region_size == 0)
return -ENODEV;
if (offset + size > region_base + region_size)
return -EINVAL;
- pfn = (vgpu->bars[2].gpa + (offset - region_base)) >> PAGE_SHIFT;
+ if (vgpu->bars[2].emulated) {
+ /* Emulation-only: map the private shadow's own pages,
+ * never the physical GPU's BAR2 PFNs. */
+ vm_flags_set(vma, VM_DONTEXPAND | VM_DONTDUMP);
+ vma->vm_page_prot =
+ pgprot_writecombine(vma->vm_page_prot);
+ for (pos = 0; pos < size; pos += PAGE_SIZE) {
+ pfn = vmalloc_to_pfn(vgpu->bars[2].vaddr +
+ bar2_off + pos);
+ if (remap_pfn_range(vma, vma->vm_start + pos,
+ pfn, PAGE_SIZE,
+ vma->vm_page_prot))
+ return -EAGAIN;
+ }
+ return 0;
+ }
+ pfn = (vgpu->bars[2].gpa + bar2_off) >> PAGE_SHIFT;
vma->vm_page_prot = pgprot_noncached(vma->vm_page_prot);
return remap_pfn_range(vma, vma->vm_start, pfn, size, vma->vm_page_prot);
}
@@ -2643,6 +2662,14 @@ module_param_named(sema_ce_grace_ms, gsp_sema_ce_grace_ms, uint, 0644);
MODULE_PARM_DESC(sema_ce_grace_ms, "CE pending-entry grace window in ms (default 8000)");
/*
+ * Bounded work-before-release hold (ms) for an unreadable CE entry; the
+ * deferred retry probe shows translations/pins land 0.5-2 s late. Absolute
+ * per-entry deadline, never reset; 0 disables.
+ */
+module_param_named(sema_push_hold_ms, gsp_sema_push_hold_ms, uint, 0644);
+MODULE_PARM_DESC(sema_push_hold_ms, "CE unreadable-entry work-before-release hold in ms (default 2500, 0 disables)");
+
+/*
* Grace window (ms) an untranslatable GR entry stays pending before the pump
* consumes it as stale, so a late push-buffer pin can still make it readable.
*/
diff --git a/tools/cuda-ab/run/vector_add b/tools/cuda-ab/run/vector_add
index c1581c1..0598212 100755
--- a/tools/cuda-ab/run/vector_add
+++ b/tools/cuda-ab/run/vector_add
Binary files differ
diff --git a/tools/cuda-ab/run/victory_mail b/tools/cuda-ab/run/victory_mail
index b04f59e..62ef1de 100755
--- a/tools/cuda-ab/run/victory_mail
+++ b/tools/cuda-ab/run/victory_mail
Binary files differ
diff --git a/tools/cuda-ab/victory_mail.cu b/tools/cuda-ab/victory_mail.cu
index e83fc1b..296f7e6 100644
--- a/tools/cuda-ab/victory_mail.cu
+++ b/tools/cuda-ab/victory_mail.cu
@@ -33,6 +33,9 @@
#include <cstdio>
#include <cstdlib>
#include <cstring>
+#include <cstdint>
+#include <fcntl.h>
+#include <unistd.h>
#include <cuda_runtime.h>
#define MAIL "the little gpu man says you've got mail"
@@ -68,6 +71,46 @@ __global__ void vadd(const float *a, const float *b, float *c, int n)
c[i] = a[i] + b[i];
}
+/* Diagnostic: where does a CUDA pointer live for the CPU?
+ * Prints the /proc/self/maps line covering it and the pagemap
+ * present/PFN entry, so the read-back mechanism of the DtoH copy is on
+ * record (CPU mapping vs device-only). Never affects control flow. */
+static void dump_va(const char *tag, const void *p)
+{
+ unsigned long long want = (unsigned long long)(uintptr_t)p;
+ char line[512];
+ FILE *m = fopen("/proc/self/maps", "r");
+
+ printf("DMAP %s va=0x%llx\n", tag, want);
+ if (m) {
+ while (fgets(line, sizeof(line), m)) {
+ unsigned long long lo, hi;
+
+ if (sscanf(line, "%llx-%llx", &lo, &hi) == 2 &&
+ want >= lo && want < hi) {
+ printf("DMAPLINE %s", line);
+ break;
+ }
+ }
+ fclose(m);
+ }
+ int pm = open("/proc/self/pagemap", O_RDONLY);
+ if (pm >= 0) {
+ unsigned long long e = 0;
+ off_t off = (off_t)((want / 4096) * 8);
+
+ if (pread(pm, &e, 8, off) == 8)
+ printf("DPRESENT %s present=%d pfn=0x%llx\n", tag,
+ (int)((e >> 63) & 1),
+ (unsigned long long)(e & ((1ULL << 55) - 1)));
+ else
+ printf("DPRESENT %s read failed\n", tag);
+ close(pm);
+ } else {
+ printf("DPRESENT %s pagemap open failed\n", tag);
+ }
+}
+
int main(void)
{
cudaDeviceProp prop;
@@ -90,17 +133,30 @@ int main(void)
/* The GPU writes the mail. */
CHECK(cudaMalloc(&d_mail, len));
+ dump_va("d_mail_alloc", d_mail);
mail_kernel<<<1, 64>>>(d_mail, len);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
memset(h_mail, 0, sizeof(h_mail));
CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost));
+ dump_va("d_mail_after_dtoh", d_mail);
if (memcmp(h_mail, MAIL, len) != 0) {
+ printf("MAILDUMP ");
+ for (i = 0; i < len; i++)
+ printf("%02x", (unsigned char)h_mail[i]);
+ printf("\n");
printf("MAIL bad=<kernel output did not match>\n");
printf("STATUS FAIL\n");
return 1;
}
printf("MAIL %s\n", h_mail);
+ {
+ size_t fm = 0, tm = 0;
+ cudaError_t me = cudaMemGetInfo(&fm, &tm);
+
+ printf("MEMINFO rc=%d free=%llu total=%llu\n", (int)me,
+ (unsigned long long)fm, (unsigned long long)tm);
+ }
/* The classic acceptance check: vectorAdd. */
h_a = (float *)malloc(bytes);
@@ -114,9 +170,27 @@ int main(void)
h_b[i] = (float)(i % 997) * 0.25f;
}
- CHECK(cudaMalloc(&d_a, bytes));
- CHECK(cudaMalloc(&d_b, bytes));
- CHECK(cudaMalloc(&d_c, bytes));
+ {
+ cudaError_t e_a = cudaMalloc(&d_a, bytes);
+ cudaError_t e_b = e_a == cudaSuccess ?
+ cudaMalloc(&d_b, bytes) : e_a;
+ cudaError_t e_c = e_b == cudaSuccess ?
+ cudaMalloc(&d_c, bytes) : e_b;
+
+ printf("MALLOC bytes=%zu d_a=%d d_b=%d d_c=%d\n", bytes,
+ (int)e_a, (int)e_b, (int)e_c);
+ if (e_a != cudaSuccess) {
+ size_t fm = 0, tm = 0;
+
+ cudaMemGetInfo(&fm, &tm);
+ printf("MEMINFO2 free=%llu total=%llu\n",
+ (unsigned long long)fm,
+ (unsigned long long)tm);
+ }
+ CHECK(e_a);
+ CHECK(e_b);
+ CHECK(e_c);
+ }
CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
diff --git a/tools/cuda-bridge/cuda_bridge b/tools/cuda-bridge/cuda_bridge
index 21fa295..d3866da 100755
--- a/tools/cuda-bridge/cuda_bridge
+++ b/tools/cuda-bridge/cuda_bridge
Binary files differ
diff --git a/tools/cuda-bridge/cuda_bridge.c b/tools/cuda-bridge/cuda_bridge.c
index 86823e6..cfd2ab7 100644
--- a/tools/cuda-bridge/cuda_bridge.c
+++ b/tools/cuda-bridge/cuda_bridge.c
@@ -1026,6 +1026,64 @@ static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in,
uint32_t bad = 0, i;
int ni;
+ /*
+ * Diagnostic (input-path proof): the acceptance client fills a with
+ * (i%1000)*0.5 and b with (i%997)*0.25. Verify the bytes this daemon
+ * actually received against that pattern, so an input-path corruption
+ * is separated from a kernel/execution corruption. Bounded output.
+ */
+ {
+ uint32_t ba = 0, bb = 0, nan_in = 0, first = 0xffffffffu, k;
+
+ for (k = 0; k < n; k++) {
+ float ea = (float)(k % 1000) * 0.5f;
+ float eb = (float)(k % 997) * 0.25f;
+
+ if (ha[k] != ea) {
+ if (first == 0xffffffffu)
+ first = k;
+ ba++;
+ }
+ if (hb[k] != eb)
+ bb++;
+ if (ha[k] != ha[k] || hb[k] != hb[k])
+ nan_in++;
+ }
+ printf("CUBB2 seq=%llu INCHECK n=%u bad_a=%u bad_b=%u nan=%u first_bad=%u a0=%08x b0=%08x a_last=%08x b_last=%08x\n",
+ (unsigned long long)j->seq, n, ba, bb, nan_in, first,
+ ((const uint32_t *)ha)[0], ((const uint32_t *)hb)[0],
+ ((const uint32_t *)ha)[n - 1], ((const uint32_t *)hb)[n - 1]);
+ fflush(stdout);
+ if (first != 0xffffffffu) {
+ printf("CUBB2 seq=%llu INSAMPLE first=%u a[first..+8]=",
+ (unsigned long long)j->seq, first);
+ for (k = first; k < first + 8 && k < n; k++)
+ printf(" %08x", ((const uint32_t *)ha)[k]);
+ printf(" b[first..+8]=");
+ for (k = first; k < first + 8 && k < n; k++)
+ printf(" %08x", ((const uint32_t *)hb)[k]);
+ printf("\n");
+ fflush(stdout);
+ }
+ /* Raw blob dump for offline analysis (bounded, overwrite). */
+ {
+ char path[64];
+ FILE *fp;
+
+ snprintf(path, sizeof(path), "/tmp/cubb2-in-seq%llu.bin",
+ (unsigned long long)j->seq);
+ fp = fopen(path, "wb");
+ if (fp) {
+ fwrite(in, 1, (size_t)bytes * 2, fp);
+ fclose(fp);
+ printf("CUBB2 seq=%llu INDMP %s bytes=%llu\n",
+ (unsigned long long)j->seq, path,
+ (unsigned long long)bytes * 2);
+ fflush(stdout);
+ }
+ }
+ }
+
hc = malloc(bytes);
if (!hc) {
d->status = ST_RAM;
@@ -1070,6 +1128,40 @@ static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in,
}
d->bad = bad;
d->checksum = double_to_bits(sum);
+ {
+ uint64_t h = 1469598103934665603ULL;
+ uint32_t k, nan_out = 0;
+
+ for (k = 0; k < n; k++) {
+ uint32_t w;
+
+ memcpy(&w, hc + (size_t)k * 4, 4);
+ h ^= w;
+ h *= 1099511628211ULL;
+ if (hf[k] != hf[k])
+ nan_out++;
+ }
+ printf("CUBB2 seq=%llu OUTCHECK n=%u hash=%016llx nan=%u bad=%u c0=%08x c_last=%08x\n",
+ (unsigned long long)j->seq, n,
+ (unsigned long long)h, nan_out, bad,
+ ((const uint32_t *)hc)[0], ((const uint32_t *)hc)[n - 1]);
+ fflush(stdout);
+ {
+ char path[64];
+ FILE *fp;
+
+ snprintf(path, sizeof(path), "/tmp/cubb2-out-seq%llu.bin",
+ (unsigned long long)j->seq);
+ fp = fopen(path, "wb");
+ if (fp) {
+ fwrite(hc, 1, (size_t)n * 4, fp);
+ fclose(fp);
+ printf("CUBB2 seq=%llu OUTDMP %s bytes=%u\n",
+ (unsigned long long)j->seq, path, n * 4);
+ fflush(stdout);
+ }
+ }
+ }
*outp = hc;
*out_len = bytes;
hc = NULL;
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 10:32:54 +0000