debug/issue: issue register + ISSUE-0001 (reproducibility) + ISSUE-0002 (multiple kernels per VM)

- debug/issue/README.md: when to file (nasty bugs only; default is no
  issues), template, numbering, lifecycle incl. solved/ move.
- issue-number-1.md: guest CUDA is not reproducibly green (top priority);
  evidence discriminators, root cause (starved pool release translation),
  the three fixes landed, 10/10-boot validation plan.
- issue-number-2.md: multiple kernels in one VM (sequential/concurrent/
  simultaneous) must work; second-client wedge root-cause chain + plan.
- tools/cuda-ab/smiley.cu: 4-pass vectorAdd that renders a 64x32 ASCII
  smiley from GPU-computed distance fields; ab.sh gains a generic 'push'
  mode so any WORKLOAD binary can be delivered to the guest.

5 files changed, 525 insertions, 0 deletions

diff --git a/debug/issue/README.md b/debug/issue/README.md
new file mode 100644
index 0000000..2c08afc
--- /dev/null
+++ b/debug/issue/README.md
@@ -0,0 +1,74 @@
+# Issue register — process
+
+**Policy: do NOT create issues by default.** An issue here is a *nasty bug*:
+something that blocks mission-critical behaviour, is non-deterministic or not
+reproducible on demand, corrupts state, or spans subsystems such that a
+single-session fix is unrealistic. Ordinary fixes, TODOs and feature work do
+not get an issue — they live in the handoff and the code comments.
+
+File an issue when **any** of these holds:
+
+1. Reproducibility / liveness is variable (works sometimes, stalls others)
+   and the variance itself is the bug.
+2. The failure blocks a 1.0 acceptance criterion (guest CUDA, multi-VM,
+   performance target, ...).
+3. Cross-subsystem state (host module + guest driver + physical GPU +
+   userspace tools) is involved and the ownership of the fix is unclear.
+4. The investigation has already burned more than one session without a
+   confirmed root cause.
+
+## How to file
+
+* Directory: `open-iov/debug/issue/`; solved ones move to
+  `open-iov/debug/issue/solved/` (never delete, never renumber).
+* Filename: `issue-number-N.md` with the next free integer `N`
+  (`issue-number-1.md`, `issue-number-2.md`, ...). Create
+  `issue-number-N+1.md` as a copy of the template below — do not create
+  placeholder issues.
+* Keep it **public-safe** (this directory is tracked): describe components,
+  behaviour and evidence (same rules as every tracked document).
+* Reference evidence by artifact: session/log file names, counter lines,
+  commit hashes. Every claim needs its evidence pointer.
+
+## Template
+
+```markdown
+# ISSUE-00NN: <one-line title>
+* Status: OPEN | FIX IMPLEMENTED (VALIDATION PENDING) | SOLVED
+* Filed: <date>, session <label> | Priority: <top/p1/p2> | Component: <kernel/guest/...>
+
+## Symptom
+What is observed, by whom, how often. Include the exact observable signature.
+
+## Impact
+Which acceptance criterion is blocked; what work is blocked behind it.
+
+## Evidence
+Logs/counters/commits that establish the facts. Include the one-line
+discriminators between healthy and failing runs.
+
+## Root cause
+The causal chain established so far, with confidence and what is still
+hypothesis (label uncertainty explicitly).
+
+## Fix / fix plan
+Code changes with commit hashes; for open items, the ordered plan and the
+files/functions to touch.
+
+## Validation plan
+Acceptance criteria (numbers of consecutive runs, counters that must stay
+zero, suites to run) and how the fix is rolled back if it regresses.
+
+## History
+Dated one-liners: what was tried, what it taught, what was rejected.
+```
+
+## Lifecycle
+
+* OPEN: filed, root cause partial.
+* FIX IMPLEMENTED (VALIDATION PENDING): code landed (committed) but the
+  acceptance runs have not completed on a live stack.
+* SOLVED: acceptance criteria met with evidence; move the file to `solved/`
+  and add a `## Resolution` section (what fixed it, the validating runs,
+  residual risks). If the same failure returns, file a **new** number and
+  cross-reference — do not reopen the moved file.
diff --git a/debug/issue/issue-number-1.md b/debug/issue/issue-number-1.md
new file mode 100644
index 0000000..f23721f
--- /dev/null
+++ b/debug/issue/issue-number-1.md
@@ -0,0 +1,128 @@
+# ISSUE-0001: Guest CUDA is not reproducibly green — the client stalls at DEVICE (GR context init never submitted)
+
+* Status: FIX IMPLEMENTED (VALIDATION PENDING)
+* Filed: 2026-10-08, session `smiley-m1`/deep-debug | Priority: **top (blocks 1.0)** | Component: kernel (`vgsp.ko` sema service + RM emulation) with a userspace-watcher race as the variance source
+
+## Symptom
+
+`victory_mail` (and every other stock-CUDA client) run in the guest does one
+of two things on a cold boot:
+
+* **Green**: prints `DEVICE …`, then `MAIL …`, `VADD n=32768 bad=0`, `STATUS OK`
+  within ~15–190 ms of engine work; host forwards CE and GR pushes
+  (`work_gr req>0`, `GUEST GR LAUNCH OK`).
+* **Stalled**: prints `DEVICE …` and then nothing. The guest RM spins in
+  `NV_ESC_RM_CONTROL` (`0x2a`, `NV2080_CTRL_CMD_MC_SERVICE_INTERRUPTS`
+  `0x20801702`) at ~1 Hz forever. The GR rings are **never written by the
+  guest**; the host sees `work_gr req=0`; FB scans find only CE ring entries.
+
+The split is per-boot. Historically (m104, m141-r3, vmkeep6/7/14) roughly
+30–50 % of boots were green on a *fresh* host; after ~8 h of host uptime and
+~20 module claim/release cycles the green rate goes to ~0 % and the failure
+can degrade further (the guest agent never comes up; smiley r1–r3).
+
+**This is not a regression against any code change: m141-r3 (the accepted
+victory, commit `895ab8b8`) was itself a lucky boot of the same intermittent
+class; the same recipe/commit stalls on other boots.**
+
+## Impact
+
+* Blocks the 1.0 acceptance: "stock guest CUDA works flawlessly, every time".
+* Forces `GRFWD_RETRIES`/`GRFWD_STALL_SECS` lottery runs (minutes per boot)
+  and hides new regressions under the stall noise.
+* Blocks everything behind it: smiley/multi-kernel validation, P2 kernels,
+  performance numbers, multi-mdev, two guests.
+
+## Evidence
+
+One-line discriminators (host dmesg counters, `sema-svc:` periodic stats):
+
+| Counter | Green (m141-r3) | Stalled (p13-r4) | Meaning |
+|---|---|---|---|
+| `pushes` / `applied` | 655 / 655 | **90 / 89, then frozen** | one completion never lands |
+| `pushes_nomap` | 0 | **1** | one release page had no translation |
+| `gr_rel` / `pub_calls` | 219 / 2801 | **0 / 0** | guest never starts GR |
+| `ring_auto` | 26 | **6** (CE only) | GR rings never created/pinned |
+| `work_gr req` | 2+ | **0** | no GR submission reaches the host |
+| `injected mapping` lines | 10 / 40 | **0** | watcher variance |
+
+Other confirmed facts:
+
+* `tools/guest-diag/guest_kprobes.sh` armed all 6 kprobes (`ok`×6) and the
+  semfix watcher translated CE sem words in stalled boots (p13-r4) — so the
+  stall is **not** the earlier "kprobes never armed" class (`6eee9ac3` fix
+  was real but is not this issue).
+* The guest's own wait is a tracker completion: wedged stacks end in
+  `uvm_tracker_wait_for_entry ← uvm_push_end_and_wait ← channel_pool_add ←
+  uvm_channel_manager_create` (repeat-client instance) — i.e. the guest is
+  waiting for a CE push completion to become visible in a semaphore/waiter
+  page.
+* `svc_apply_pending()` keeps a release that cannot translate ("not
+  translatable yet: keep") and retries every pass; nothing else ever moves it.
+* Related (separate correctness bug, fixed under this issue): the
+  `PROMOTE_CTX` (0x2080012b) reply. See "Fixes".
+
+## Root cause
+
+The completion (release) of the client's last CE push in the boot-critical
+window cannot be written back because its semaphore/waiter page has no
+VA→PA translation yet. The release is deferred into `svc_pending[]` and waits
+for a translation that arrives only from the userspace watcher pipeline
+(`pin_gr_stream`/`semfix_uvmt`) or from the kernel's *partial* fallback set.
+The watcher pipeline is racy across boots (kprobe arming vs the client's
+first push, scanner vs pool allocation, kcore scans vs tracker windows), so
+on some boots every fallback misses; the guest's tracker wait never wakes;
+the client never creates its GR channels; nothing ever reaches the GR rings.
+
+Kernel fallback coverage was incomplete exactly for the pool/waiter page: the
+pool resolver (`svc_pool_pa_resolve`) had no walk of the **driver-published
+VAS page directories** (`svc_walk_vaspaces`) — the deterministic, watcher-free
+source the general resolver (`svc_lookup_trusted`) has always used.
+
+Confidence: high for the mechanism (the deferral + the missing walk + the
+observed freeze at `applied=89/90`). The exact *reason the walk itself misses
+on some boots* (directory registered later than the first push, or the
+watcher path succeeding first on greens) remains boot-timing dependent; the
+fix below removes the dependence rather than chasing the timing.
+
+## Fix
+
+1. **`kernel/gsp_sema_service.c` — `svc_pool_pa_resolve`:** add
+   `svc_walk_vaspaces()` (driver-published directories, read-back proven)
+   between the self-discovered walk and the targeted rm_mem scan. Kernel-native
+   translation of the deferred release page no longer depends on the watcher.
+2. **`kernel/gsp_sema_service.c` — `svc_pending` starvation telemetry:**
+   per-entry `tries` counter + bounded logs at 4/64/1024 attempts naming
+   `va`, `payload` and the consequence. A future starvation is visible in
+   dmesg instead of silent.
+3. **`kernel/gsp_rm_emulation.c` — PROMOTE_CTX reply (supersedes
+   `fa9296af`):** the reply must keep `bInitialize` set. Per
+   `kernel_graphics_object.c` `kgrobjPromoteContext_IMPL`:
+   `if (!params.promoteEntry[i].bInitialize) continue;
+   kgrctxMarkCtxBufferInitialized(...)` — the flag is the server's
+   "initialized" ACK; clearing it (fa9296af) made every ctx buffer re-promote
+   forever. This was a correctness bug in the stall's vicinity, not the stall
+   itself (the stall predates and outlives it).
+
+Durable direction (if a stall variant survives the fix): fold the watcher's
+job into the kernel entirely (complete `svc_pb_discover`/`svc_pd` coverage of
+the CE push/pool windows) so the watchers are accelerators only.
+
+## Validation plan
+
+1. Fresh host boot (reboot approved 2026-10-08): accepted recipe, plain
+   runner — expect `victory_mail` green.
+2. **10/10 consecutive cold boots** green within seconds; zero stall
+   captures; `pending release starved` never logs; `pushes_nomap` returns 0
+   on every green.
+3. Re-run the smiley workload (`tools/cuda-ab/smiley.cu`) per green VM.
+4. Rollback: revert commits 1–3 independently; rebuild; the previous
+   behaviour (lottery) returns without new failure modes.
+
+## History
+
+* 2026-10-02..07: stall hunted through m1xx series; greens m104, m141-r3;
+  kprobe arming fixed (`6eee9ac3`); plan to kernel-native discovery recorded.
+* 2026-10-08: campaigns p1–p13 + smiley r1–r3 all stall; evidence shows the
+  missing completion + missing pool-page translation; `fa9296af` PROMOTE_CTX
+  change tested and disproven (and is itself corrected here).
diff --git a/debug/issue/issue-number-2.md b/debug/issue/issue-number-2.md
new file mode 100644
index 0000000..e04e08f
--- /dev/null
+++ b/debug/issue/issue-number-2.md
@@ -0,0 +1,92 @@
+# ISSUE-0002: Multiple CUDA kernels/workloads inside one VM — sequential, concurrent and simultaneous must all work
+
+* Status: OPEN (root cause established for the sequential case; concurrent/simultaneous untested)
+* Filed: 2026-10-08 | Priority: **p1 (blocks 1.0 "multiple kernels", follows ISSUE-0001)** | Component: kernel per-client/session state + guest UVM/RM interaction
+
+## Symptom
+
+* **Sequential (known):** a *second* CUDA process in the same boot wedges:
+  spins at ~100 % CPU with no output (not even `DEVICE`), host traffic for it
+  is zero. The first client in the same boot is green (10–34 s). Until fixed,
+  operations are "one client per boot, reboot for a rerun".
+* **Concurrent:** not yet tested (blocked by the sequential wedge).
+* **Simultaneous (one process, multiple streams/kernels at once):** not yet
+  tested; expected to hit the same per-channel/per-client state class.
+
+## Impact
+
+* Blocks 1.0 acceptance "arbitrary CUDA kernels flawlessly" beyond a single
+  sequential run, all realistic workloads (frameworks spawn many processes),
+  and the P2 kernel suite (`tools/cuda-ab/p2/run_guest.sh`) which runs
+  multiple programs per boot by design.
+
+## Evidence
+
+* Wedged guest stack (vmkeep10 run C, pid 1447):
+  `uvm_ioctl → uvm_api_register_gpu → uvm_va_space_register_gpu →
+  uvm_gpu_retain_by_uuid+0x44 [nvidia_uvm]` (main thread), second thread in
+  `do_poll`.
+* Source (guest UVM, `kernel-open/nvidia-uvm/uvm_gpu.c:3129-3163`):
+  `uvm_gpu_retain_by_uuid` takes `g_uvm_global.global_lock`;
+  `uvm_gpu_release_locked` (last-release path) holds the same lock while
+  calling `nvUvmInterfaceUnregisterGpu` — an RM call to the host. If that
+  exchange never completes host-side, the lock is held forever and every
+  later client wedges at registration.
+* Host view (vmkeep10): run A full RM traffic (channel `0xcaf000d4`
+  registered), then idle; run B issues **zero** RM/sema traffic; each client
+  run re-runs the GSP boot handshake (`GSP_INIT_DONE` per run start), so the
+  host must fully re-arm per session. Around run A's exit:
+  `sema-svc: completion hint rc=-22` twice on the fresh channel.
+* Host counters freeze for the second run (`counter pub_calls=2790
+  gr_rel=216`, `drain held=2`, `faked=0`).
+* Evidence pack: `backups/evidence-20261002-victory/analysis-2026-10-08-repeat-stall/`
+  (`run1/2/3.log`, `dmesg-w-vmkeep10.log`, semfix/pin logs).
+* An m146/m147 experiment (re-arm per-rank one-shot latches on channel
+  registration; derive CE rank from ring VA; re-arm on GSP session reset via
+  `gsp_sema_service_session_reset()` from `gsp_rm_emu_reset`) was implemented
+  and **reverted** — it did not fix run B and temporarily stopped first-run
+  greens from reproducing. WIP diff:
+  `…/analysis-2026-10-08-repeat-stall/m146-m147-session-reset-WIP.patch`.
+
+## Root cause
+
+Established for the sequential case: per-client state in the host sema/GR
+service is not reset when a new guest client context appears (the guest
+re-runs the GSP boot handshake per client). The old client's channel/rank/
+ring/pool state makes the host's completion path go quiet for the new client
+(`faked=0`, `drain held=2`, zero new applies), so the new client's UVM
+registration RM exchange never completes — and because the guest UVM holds
+`global_lock` across that exchange (`uvm_gpu_release_locked` →
+`nvUvmInterfaceUnregisterGpu`), the whole guest UVM wedges for every later
+client too. Concurrent/simultaneous cases are expected to share this class
+(per-channel one-shot latches, single per-vgpu sema context).
+
+## Fix / fix plan
+
+1. Reset/rebind host per-client/per-channel sema+GR state when a new client
+   context appears (new GSP session / new channel registration): re-arm
+   one-shot rank latches, clear stale ring/userd publications, bound any
+   drain hold so a miss degrades instead of freezing (the `drain held` path).
+2. Per-vgpu sema context + P0 isolation work: ordered hunks in
+   `backups/evidence-20261002-victory/analysis-2026-10-07/sa-p2-patchnotes.md`
+   (P0-1..P0-7, stages A1–A6, gates G1/G3, `users/idle` anti-UAF drain).
+3. Graduation criteria: sequential 3+ clients in one boot green; two
+   concurrent processes (different streams) green; one process with N
+   simultaneous streams/kernels green; then the P2 suite in a single boot.
+
+## Validation plan
+
+* Sequential: 3 consecutive `victory_mail` runs in one VM, all green.
+* Concurrent: two `victory_mail` instances launched together, both green.
+* Simultaneous: multi-stream kernel test (P2 suite addition) green,
+  oracle-matched against bare metal.
+* Rollback: patch reverts independently (see WIP diff for the reverted
+  experiment).
+
+## History
+
+* 2026-10-08 00:05: first-client greens became deterministic (`6eee9ac3`);
+  second-client wedge discovered the same night (vmkeep9–14).
+* 2026-10-08 01:00: root-cause chain assembled; m146/m147 experiment
+  reverted as ineffective; kernel-native per-client reset identified as the
+  durable path.
diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh
index 8ae05ee..5a290e4 100755
--- a/tools/cuda-ab/ab.sh
+++ b/tools/cuda-ab/ab.sh
@@ -86,6 +86,7 @@ host)    run_host ;;
 guest)   run_guest ;;
 compare) compare ;;
 all)     run_host; run_guest; compare ;;
+push)    push_bin ;;
 victory)
 	WORKLOAD=victory_mail
 	SRC="$HERE/$WORKLOAD.cu"
diff --git a/tools/cuda-ab/smiley.cu b/tools/cuda-ab/smiley.cu
new file mode 100644
index 0000000..42770cc
--- /dev/null
+++ b/tools/cuda-ab/smiley.cu
@@ -0,0 +1,230 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * smiley.cu - vectorAdd draws a smiley face.
+ *
+ * Four vectorAdd passes execute on the GPU; each computes one scalar field
+ *
+ *   c[i] = a[i] + b[i]
+ *
+ * over the same n = 32768 elements laid out as a 256x128 sample grid:
+ *
+ *   pass 0  a=u^2          b=v^2          -> face disk field
+ *   pass 1  a=(u+xe)^2     b=(v+ye)^2     -> left-eye disk field
+ *   pass 2  a=(u-xe)^2     b=(v+ye)^2     -> right-eye disk field
+ *   pass 3  a=u^2          b=(v-ym)^2     -> mouth annulus field
+ *
+ * with u,v centre-relative, aspect-corrected coordinates. The host copies
+ * the four fields back and thresholds them to render a 64x32 ASCII smiley
+ * (4x4 samples per character cell). The artwork therefore comes out of GPU
+ * addition results only - no GPU-side drawing, no CPU-side math on pixels
+ * beyond the thresholds.
+ *
+ * Output lines are stable and greppable (matches tools/cuda-ab/ab.sh):
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   MALLOC bytes=<n> rc=<0|err>
+ *   VADD passes=4 n=<count> bad=<mismatches> checksum=<sum>
+ *   RESULT cells=<W>x<H> inked=<cells> checksum=<sum>
+ *   STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation:
+ *
+ *   nvcc -O2 --cudart=static -o smiley smiley.cu
+ *
+ * Then, from the guest:  /tmp/smiley
+ */
+
+#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];
+}
+
+/* sample grid */
+#define GW 256
+#define GH 128
+#define GN (GW * GH)
+
+/* character cell = 4x4 samples */
+#define CW 64
+#define CH 32
+#define CELL 4
+
+/* feature geometry, in aspect-corrected grid units (u ~ +/-127, v ~ +/-127) */
+#define FACE_R2   10000.0f /* face disk:    u^2+v^2   <= 100^2   */
+#define EYE_X     36.0f
+#define EYE_Y     30.0f
+#define EYE_R2    144.0f   /* eyes: r=12 */
+#define MOUTH_Y   24.0f
+#define MOUTH_IN  2025.0f  /* 45^2 */
+#define MOUTH_OUT 3025.0f  /* 55^2 */
+#define MOUTH_V   40.0f
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	float *h_a, *h_b, *h_out, *fields;
+	float *d_a, *d_b, *d_c;
+	size_t bytes = (size_t)GN * sizeof(float);
+	double checksum = 0.0;
+	int bad = 0, pass, i;
+
+	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);
+
+	{
+		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 rc=%d\n", bytes, (int)e_a);
+		if (e_a != cudaSuccess) {
+			size_t fm = 0, tm = 0;
+
+			cudaMemGetInfo(&fm, &tm);
+			printf("MEMINFO free=%llu total=%llu\n",
+			       (unsigned long long)fm,
+			       (unsigned long long)tm);
+		}
+		CHECK(e_a);
+		CHECK(e_b);
+		CHECK(e_c);
+	}
+
+	h_a = (float *)malloc(bytes);
+	h_b = (float *)malloc(bytes);
+	h_out = (float *)malloc(bytes);
+	fields = (float *)malloc((size_t)4 * bytes);
+	if (!h_a || !h_b || !h_out || !fields)
+		return 1;
+
+	for (pass = 0; pass < 4; pass++) {
+		for (i = 0; i < GN; i++) {
+			int x = i % GW, y = i / GW;
+			float u = (float)x - 127.5f;
+			float v = ((float)y - 63.5f) * 2.0f;
+			float pa, pb;
+
+			switch (pass) {
+			case 0: /* face */
+				pa = u * u;            pb = v * v;
+				break;
+			case 1: /* left eye */
+				pa = (u + EYE_X) * (u + EYE_X);
+				pb = (v + EYE_Y) * (v + EYE_Y);
+				break;
+			case 2: /* right eye */
+				pa = (u - EYE_X) * (u - EYE_X);
+				pb = (v + EYE_Y) * (v + EYE_Y);
+				break;
+			default: /* mouth */
+				pa = u * u;
+				pb = (v - MOUTH_Y) * (v - MOUTH_Y);
+				break;
+			}
+			h_a[i] = pa;
+			h_b[i] = pb;
+		}
+
+		CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
+		CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
+		vadd<<<(GN + 255) / 256, 256>>>(d_a, d_b, d_c, GN);
+		CHECK(cudaGetLastError());
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_out, d_c, bytes, cudaMemcpyDeviceToHost));
+
+		for (i = 0; i < GN; i++) {
+			if (h_out[i] != h_a[i] + h_b[i])
+				bad++;
+			checksum += (double)h_out[i];
+		}
+		for (i = 0; i < GN; i++)
+			fields[(size_t)pass * GN + i] = h_out[i];
+	}
+
+	printf("VADD passes=4 n=%d bad=%d checksum=%.0f\n", GN, bad, checksum);
+	if (bad != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	{
+		static char art[CH][CW + 1];
+		int inked = 0, cy, cx;
+
+		for (cy = 0; cy < CH; cy++) {
+			for (cx = 0; cx < CW; cx++) {
+				int hit = 0, sy, sx;
+				const float *f0 = fields;
+				const float *f1 = fields + GN;
+				const float *f2 = fields + 2 * GN;
+				const float *f3 = fields + 3 * GN;
+
+				for (sy = 0; sy < CELL; sy++) {
+					int y = cy * CELL + sy;
+					float v = ((float)y - 63.5f) * 2.0f;
+
+					for (sx = 0; sx < CELL; sx++) {
+						int x = cx * CELL + sx;
+						int idx = y * GW + x;
+						int inside;
+
+						inside = f0[idx] <= FACE_R2;
+						if (f1[idx] <= EYE_R2 ||
+						    f2[idx] <= EYE_R2)
+							inside = 0;
+						if (f3[idx] >= MOUTH_IN &&
+						    f3[idx] <= MOUTH_OUT &&
+						    v >= MOUTH_V)
+							inside = 0;
+						hit += inside;
+					}
+				}
+				art[cy][cx] = hit >= CELL * CELL / 2 ? '#' :
+									 ' ';
+				inked += art[cy][cx] == '#';
+			}
+			art[cy][CW] = 0;
+		}
+
+		printf("RESULT cells=%dx%d inked=%d checksum=%.0f\n",
+		       CW, CH, inked, checksum);
+		for (cy = 0; cy < CH; cy++)
+			printf("%s\n", art[cy]);
+	}
+
+	printf("STATUS OK\n");
+	CHECK(cudaFree(d_a));
+	CHECK(cudaFree(d_b));
+	CHECK(cudaFree(d_c));
+	free(h_a);
+	free(h_b);
+	free(h_out);
+	free(fields);
+	return 0;
+}