diff options
| -rw-r--r-- | kernel/gsp_sema_service.c | 141 | ||||
| -rw-r--r-- | kernel/gsp_sema_service.h | 3 | ||||
| -rw-r--r-- | kernel/main.c | 2 | ||||
| -rwxr-xr-x | tools/cuda-ab/run/streams2d | bin | 0 -> 829296 bytes | |||
| -rw-r--r-- | tools/cuda-ab/streams2d.cu | 197 |
5 files changed, 327 insertions, 16 deletions
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index f803309..add8b29 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -4048,7 +4048,19 @@ static u32 svc_d2hmap_log; /* D2H-MAP per-page destination budget */ * 5 ACQ_NOR = ~(v | p) != 0. */ #define SVC_WAIT_MAX 8 -#define SVC_WAIT_HOLD_MS 100u /* bounded acquire wait (fail-open) */ +/* + * r401: bounded acquire wait (fail-open), now a module parameter. The + * original 100 ms bound was shorter than a single GR launch's completion + * latency in this emulation (observed 0.5-1.6 s per launch), so a copy + * that waited on a launch's tracker (streams2: the D2H readback's + * ACQ_CIRC_GEQ on the scale channel's progress counter) gave up and ran + * BEFORE the launch wrote its output - the host read pre-scale data + * (bad=32767, c2 all wrong). 4 s covers several launch latencies; the + * bound still guarantees the entry is never parked forever (a genuinely + * unsatisfiable acquire fails open late, after the dependency's own work + * would have landed - late-but-correct instead of early-and-wrong). + */ +unsigned gsp_sema_wait_hold_ms = 4000; struct svc_wait_op { u64 va; u32 payload; @@ -5285,6 +5297,82 @@ static u32 svc_gr_cand_pos; static u64 svc_gr_cap_hist[SVC_GR_CAP_HIST]; static u32 svc_gr_cap_hist_pos; +/* + * r401: robust session-scoped "already captured" memory for the r383 + * replay-capture path. The 4-slot candidate ring + 8-VA history evicted a + * real capture within ~12 later captures, so a true replay of an + * already-forwarded launch (identical push VA and content) was + * re-captured and re-executed: the r399 soak log shows the same nine + * launch pushes re-submitted every ~10-11 s for minutes (idempotent + * kernels survive it; accumulators would not). The open-addressed table + * keeps the (push_va, push_hash) identity of every capture of the + * session, so a true replay always classifies as "seen" (LAUNCH-SKIP, + * exactly once) while a genuinely new push - even at the same VA with + * the same entry words - is captured through the normal + * `trusted && !replay` path and never consults this table. Cleared by + * svc_session_flush (a new client's identical-payload pushes restart + * their release counters, so a boot-scoped table would wrongly suppress + * them). Service-thread only. + */ +#define SVC_GR_SEEN_SLOTS 2048 +#define SVC_GR_SEEN_PROBE 8 +static struct { + u64 push_va; + u64 hash; +} svc_gr_seen[SVC_GR_SEEN_SLOTS]; +static u64 svc_gr_seen_stale; /* probe-chain evictions (telemetry) */ + +static u32 svc_gr_seen_mix(u64 push_va, u64 hash) +{ + u64 k = push_va * 0x9E3779B97F4A7C15ULL; + + k ^= hash + 0x9E3779B97F4A7C15ULL + (k << 6) + (k >> 2); + return (u32)(k % SVC_GR_SEEN_SLOTS); +} + +static bool svc_gr_seen_lookup(u64 push_va, u64 hash) +{ + u32 i, base = svc_gr_seen_mix(push_va, hash); + + if (!push_va) + return false; + for (i = 0; i < SVC_GR_SEEN_PROBE; i++) { + u32 k = (base + i) % SVC_GR_SEEN_SLOTS; + + if (!svc_gr_seen[k].push_va) + return false; + if (svc_gr_seen[k].push_va == push_va && + svc_gr_seen[k].hash == hash) + return true; + } + return false; +} + +static void svc_gr_seen_store(u64 push_va, u64 hash) +{ + u32 i, base = svc_gr_seen_mix(push_va, hash); + + if (!push_va) + return; + for (i = 0; i < SVC_GR_SEEN_PROBE; i++) { + u32 k = (base + i) % SVC_GR_SEEN_SLOTS; + + if (!svc_gr_seen[k].push_va || + (svc_gr_seen[k].push_va == push_va && + svc_gr_seen[k].hash == hash)) { + svc_gr_seen[k].push_va = push_va; + svc_gr_seen[k].hash = hash; + return; + } + } + /* Full probe chain: replace the tail (oldest-ish) entry. */ + svc_gr_seen_stale++; + svc_gr_seen[(base + SVC_GR_SEEN_PROBE - 1) % SVC_GR_SEEN_SLOTS].push_va = + push_va; + svc_gr_seen[(base + SVC_GR_SEEN_PROBE - 1) % SVC_GR_SEEN_SLOTS].hash = + hash; +} + static_assert(SVC_GR_QMD_WORDS == GSP_PHYS_GR_QMD_WORDS, "the sema GR QMD word count must match the backend ABI"); static_assert(SVC_GR_PUSH_WORDS == GSP_PHYS_GR_RAW_WORDS, @@ -5399,6 +5487,9 @@ static void svc_gr_cand_capture(struct svc_channel *ch, svc_gr_cap_hist_pos++; if (push && push_len) svc_gr_cand.push_hash = svc_push_hash(push, push_len); + /* r401: the session-wide exactly-once memory (survives ring rotation; + * see svc_gr_seen_lookup). */ + svc_gr_seen_store(push_va, svc_gr_cand.push_hash); if (push && push_len >= 4) { u32 total = push_len / 4; u32 n = total < SVC_GR_PUSH_WORDS ? total : SVC_GR_PUSH_WORDS; @@ -7052,12 +7143,12 @@ static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops, * 3. drain the pending-translation release queue (releases OUR * machinery is merely late about) and re-read once; * 4. still unsatisfied -> hold the entry (release withheld through - * the op-pending machinery) for at most SVC_WAIT_HOLD_MS per - * episode, then fail open loudly ('give-up': the value cannot - * be installed by firmware alone - it sits behind an upstream - * opdef/operand hold - so the old pass-through behaviour is - * safer than parking the channel forever). Unreadable - * semaphore: fail-open (the pre-r336 behaviour). + * the op-pending machinery) for at most gsp_sema_wait_hold_ms + * per episode, then fail open loudly ('give-up': the value + * cannot be installed by firmware alone - it sits behind an + * upstream opdef/operand hold - so the old pass-through + * behaviour is safer than parking the channel forever). + * Unreadable semaphore: fail-open (the pre-r336 behaviour). * The pass/hold/give-up decisions are logged as 'sema-wait' rows * (op, payload, observed value, decision). */ @@ -7116,7 +7207,7 @@ static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops, continue; } pend = svc_wait_pend_release(w->va); - if (waited_ms >= SVC_WAIT_HOLD_MS) { + if (waited_ms >= gsp_sema_wait_hold_ms) { /* * Bounded episode expired: fail open for the * WHOLE entry (one row, then the ops run). A @@ -8388,12 +8479,16 @@ static u64 svc_gr_nocap_log; /* * r383: has a launch with this push VA ever been captured? Scans the live - * candidate ring and the recent-capture history. Used by the replayed-QMD - * dispatch below: a replay whose push was never captured is a launch that - * would otherwise be silently dropped (the streams2 second-vadd class); a - * replay of a captured/offered push is the exactly-once no-op. + * candidate ring, the recent-capture history and the r401 session-wide + * (push_va, hash) seen table. Used by the replayed-QMD dispatch below: a + * replay whose push was never captured is a launch that would otherwise be + * silently dropped (the streams2 second-vadd class); a replay of a + * captured/offered push is the exactly-once no-op. The hash is required + * for the seen table so a later, DIFFERENT push at the same VA (a loop + * iteration: same entry words, new content) is never misclassified as + * seen - the caller only consults this on a true replay anyway. */ -static bool svc_gr_cand_push_seen(u64 push_va) +static bool svc_gr_cand_push_seen(u64 push_va, u64 hash) { u32 i; @@ -8406,7 +8501,7 @@ static bool svc_gr_cand_push_seen(u64 push_va) for (i = 0; i < SVC_GR_CAP_HIST; i++) if (svc_gr_cap_hist[i] == push_va) return true; - return false; + return svc_gr_seen_lookup(push_va, hash); } /* @@ -8543,11 +8638,15 @@ static void svc_qmd_accum(struct nv_vgpu_device *vgpu, struct svc_channel *ch, n, dup); } if (gsp_phys_gr_forward_enabled()) { + u64 push_hash = (push && push_len) ? + svc_push_hash(push, push_len) : 0; + if (trusted && !replay) { svc_gr_cand_capture(ch, st, recs, n, push, push_len, st->push_va); } else if (trusted && replay && - !svc_gr_cand_push_seen(st->push_va)) { + !svc_gr_cand_push_seen(st->push_va, + push_hash)) { /* * r383 (streams2 second-vadd class): a * complete QMD first seen on a replay pass @@ -8570,7 +8669,8 @@ static void svc_qmd_accum(struct nv_vgpu_device *vgpu, struct svc_channel *ch, svc_gr_nocap_log++; pr_info("sema-gr: LAUNCH-SKIP h=0x%x push=0x%llx (replay; already captured=%d trusted=%d)\n", ch->h_object, st->push_va, - svc_gr_cand_push_seen(st->push_va), + svc_gr_cand_push_seen(st->push_va, + push_hash), trusted); } else if (!replay) { gsp_cuda_bridge_launch(vgpu, ch->h_object, @@ -10958,6 +11058,15 @@ static void svc_session_flush(void) svc_hold_fb_log = 0; svc_cand_ovw_log = 0; svc_inline_recs_log = 0; + /* + * r401: the exactly-once capture memory is session-scoped. A new + * client's pushes restart their release counters, so their content + * (and hence hash) can collide with the previous client's - keeping + * the table would wrongly classify a fresh launch of the new + * session as already captured (LAUNCH-SKIP, silent drop). + */ + memset(svc_gr_seen, 0, sizeof(svc_gr_seen)); + svc_gr_seen_stale = 0; /* m176: new session generation - channels and scrubbers re-stamp on * registration and never inherit per-session learned state. */ svc_sess_gen++; diff --git a/kernel/gsp_sema_service.h b/kernel/gsp_sema_service.h index ea025ca..3765cd2 100644 --- a/kernel/gsp_sema_service.h +++ b/kernel/gsp_sema_service.h @@ -310,6 +310,9 @@ extern bool gsp_sema_pool_hold_fresh; */ extern bool gsp_sema_copy_dst_pin_guard; extern bool gsp_sema_wait_gate; +/* r401: bounded acquire-wait hold (fail-open) in ms; must exceed the + * emulated launch completion latency for copy-after-launch ordering. */ +extern unsigned gsp_sema_wait_hold_ms; /* * r384 (guest boot-corruption lottery): fill-destination guards (default diff --git a/kernel/main.c b/kernel/main.c index 92754d1..d2a44d6 100644 --- a/kernel/main.c +++ b/kernel/main.c @@ -2829,6 +2829,8 @@ MODULE_PARM_DESC(sema_pool_hold_fresh, "apply pool releases at fresh-pin time in */ module_param_named(sema_copy_dst_pin_guard, gsp_sema_copy_dst_pin_guard, bool, 0644); module_param_named(sema_wait_gate, gsp_sema_wait_gate, bool, 0644); +module_param_named(sema_wait_hold_ms, gsp_sema_wait_hold_ms, uint, 0644); +MODULE_PARM_DESC(sema_wait_hold_ms, "bounded acquire-wait hold before fail-open, ms (r401, default 4000; must exceed the emulated launch latency)"); MODULE_PARM_DESC(sema_copy_dst_pin_guard, "hold device-dst copies without a pin-resolved translation (default 1)"); /* diff --git a/tools/cuda-ab/run/streams2d b/tools/cuda-ab/run/streams2d Binary files differnew file mode 100755 index 0000000..90b6df8 --- /dev/null +++ b/tools/cuda-ab/run/streams2d diff --git a/tools/cuda-ab/streams2d.cu b/tools/cuda-ab/streams2d.cu new file mode 100644 index 0000000..04a0866 --- /dev/null +++ b/tools/cuda-ab/streams2d.cu @@ -0,0 +1,197 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2d.cu - streams2 with per-array / per-range diagnostic reporting. + * + * Same workload as streams2.cu (two streams, cross-stream event ordering) + * but instead of one "bad" count it reports, per array and per 4 KB block, + * how many elements mismatched, and prints the device pointers so a host + * log can be correlated with what the engine read/wrote. + * + * Output lines: + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * PTR <name>=0x<hex> (d_a1 d_b1 d_c1 d_a2 d_b2 d_c2) + * DIFF c1 total=<n> blocks=<first block index with mismatches> + * DIFF c2 total=<n> blocks=<...> + * STREAMS n=<count> bad=<mismatches> + * TIMING ms=<elapsed> + * STATUS OK|FAIL + * + * Build: nvcc -O2 --cudart=static -o streams2d streams2d.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cuda_runtime.h> + +#define CHECK(x) \ + do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("CUDA_ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("STATUS FAIL\n"); \ + return 1; \ + } \ + } while (0) + +__global__ void vadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +static void report_diff(const char *name, const float *got, + const float *exp, int n) +{ + int i, bad = 0, first = -1, last = -1; + int blk_hits = 0, blk_first = -1; + int nb = (n + 1023) / 1024; + int b; + + for (i = 0; i < n; i++) { + if (got[i] != exp[i]) { + bad++; + if (first < 0) + first = i; + last = i; + } + } + for (b = 0; b < nb; b++) { + int j, hit = 0; + + for (j = b * 1024; j < (b + 1) * 1024 && j < n; j++) + if (got[j] != exp[j]) { + hit = 1; + break; + } + if (hit) { + blk_hits++; + if (blk_first < 0) + blk_first = b; + } + } + printf("DIFF %s total=%d first=%d last=%d blocks=%d/%d blk_first=%d\n", + name, bad, first, last, blk_hits, nb, blk_first); + if (first >= 0) + printf("DIFF %s sample got[%d]=%a exp[%d]=%a got[%d]=%a exp[%d]=%a\n", + name, first, got[first], first, exp[first], last, + got[last], last, exp[last]); +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + int bad = 0, i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem); + + h_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + if (!h_a1 || !h_b1 || !h_c1 || !h_a2 || !h_b2 || !h_c2) + return 1; + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + printf("PTR d_a1=%p d_b1=%p d_c1=%p d_a2=%p d_b2=%p d_c2=%p\n", + (void *)d_a1, (void *)d_b1, (void *)d_c1, + (void *)d_a2, (void *)d_b2, (void *)d_c2); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + report_diff("c1", h_c1, h_a1, n); + { + float *exp2 = (float *)malloc(bytes); + + for (i = 0; i < n; i++) + exp2[i] = (h_a2[i] + h_b2[i]) * s2; + report_diff("c2", h_c2, exp2, n); + for (i = 0; i < n; i++) { + if (h_c1[i] != h_a1[i] + h_b1[i]) + bad++; + if (h_c2[i] != exp2[i]) + bad++; + } + free(exp2); + } + + printf("STREAMS n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + return 0; +} |