The SnapDeals OOM: A Crash Report That Exposed a Deeper Architectural Flaw

Introduction

In the middle of an intense debugging session focused on multi-GPU proof generation for Filecoin's proving infrastructure, a new crash report landed. The user pasted a wall of journalctl logs from a remote host named p-dev-ngw-1, showing the CuZK proving engine daemon crashing with an out-of-memory (OOM) error during SnapDeals proof generation. The message was brief in its framing—"New issue"—but the log it contained told a complex story of a system under strain, revealing not just a resource exhaustion problem but a fundamental architectural mismatch between the proving engine's GPU routing logic and the memory requirements of the SnapDeals circuit.

This article examines that single message in depth: why it was written, what it reveals about the system, the assumptions embedded in its delivery, and the critical knowledge it contributed to the ongoing debugging effort. The message sits at a pivotal moment in the conversation, arriving just as the assistant had diagnosed one class of GPU bug and was implementing a fix, only to be confronted with a second, seemingly distinct failure mode on different hardware.

The Message: A Crash in Progress

The user's message consists almost entirely of a raw log dump from journalctl, spanning approximately two minutes of real time on March 6, 2026. The host is p-dev-ngw-1, a machine equipped with an NVIDIA RTX 4000 Ada GPU with 20 GB of VRAM. The log traces the lifecycle of a single SnapDeals proving job from initialization through synthesis and into the GPU proving phase, where it terminates in a panic:

thread 'tokio-runtime-worker' panicked at /tmp/czk/extern/supraseal-c2/src/lib.rs:311:9:
cudaMallocAsync(&d_ptr, sz, stream)@sppark-0.1.14/sppark/util/gpu_t.cuh:73 failed: "out of memory"

The process then crashes with a segmentation fault (signal 11) and systemd records a core dump. The final line, appended by the user, provides the essential context: "On a rtx4000ada, 20GB, WindowPoSt works fine but Snap seems to OOM vram immediately post-synth."

This is not a bug report in the traditional sense—there is no structured description, no expected-versus-actual behavior, no steps to reproduce. It is a raw diagnostic dump, the kind of message a developer sends when they want to say "look at this" without filtering or interpretation. The user trusts that the assistant (and the reader of this article) can parse the log and extract the meaningful signal from the noise.

Why This Message Was Written: The Motivation and Context

To understand why this message was written, we must understand the state of the conversation at the moment it arrived. The preceding messages (indices 377–429) document a prolonged debugging effort focused on a multi-GPU data race in the PoRep (Proof of Replication) proving pipeline. The assistant had identified that the C++ GPU proving code always routed single-circuit proofs to GPU 0, regardless of which Rust worker submitted them. On a two-GPU system, this meant that workers assigned to GPU 1 would still execute their CUDA kernels on GPU 0, causing data corruption when multiple workers ran concurrently. The initial "fix" was a shared mutex that serialized all partitioned proofs onto GPU 0, effectively wasting the second GPU.

At the time of message 430, the assistant was in the process of implementing a proper solution: threading a gpu_index parameter through the entire call chain—from the Rust engine's GPU worker code, through the bellperson prover functions, across the Rust FFI boundary, and into the C++ groth16_cuda.cu code—so that the C++ layer would use the GPU assigned by the Rust engine instead of always defaulting to GPU 0. This was a multi-file, multi-language refactoring effort that touched the C++ CUDA code, the Rust FFI bindings, the bellperson integration layer, and the pipeline engine.

The user's message interrupts this work. It reports a new issue on a different host (p-dev-ngw-1 instead of cs-calib), with a different proof type (SnapDeals instead of PoRep), and a different symptom (OOM instead of proof corruption). The timing is significant: the assistant had not yet deployed the multi-GPU fix anywhere, and the shared mutex hack was still the active code on all hosts. The crash on p-dev-ngw-1 is happening under that hack.

The user's motivation is clear: they are running SnapDeals workloads on a 20 GB GPU and seeing crashes. They need the proving pipeline to work for SnapDeals, and this crash is a blocker. By sending the raw log, they are asking the assistant to diagnose the failure and determine whether it is a new bug, a manifestation of the same underlying issue, or a resource constraint that requires a different approach.

Anatomy of the Crash: A Timeline

The log spans approximately 83 seconds from SRS loading to the OOM panic. Let us trace the sequence of events:

T+0s (03:18:53): The SRS (Structured Reference String) for the snap-32g circuit loads successfully in 23 seconds. The engine dispatches a SnapDeals synthesis job with 16 partitions and 5 partition workers. Background PCE (Pre-Compiled Circuit Encoding) extraction begins.

T+0s to T+83s: Partitions 0 through 4 begin synthesis using the standard path (synthesize_with_hint). The logs repeatedly note "no capacity hint cached — first synthesis will grow organically," indicating this is the first time this circuit has been synthesized on this host. The PCE extraction runs in the background, recording the constraint structure of the circuit.

T+83s (03:20:16): PCE extraction completes. The extracted circuit is enormous: 5 inputs, 80,978,488 auxiliary variables, 81,049,504 constraints, with a total of 443,680,531 non-zero entries across the A, B, and C matrices. The summary reports "mem: 15.8 GiB." This is a critical datum: the pre-compiled circuit representation alone requires 15.8 GB of memory. The PCE fails to save to disk (non-fatal) because it cannot create a temp file in /data/zk/params/.

T+83s to T+84s: Partition 2 completes synthesis in 83.9 seconds and is sent to the GPU. Worker 1 (assigned to GPU 1) picks it up. The log shows CUZK_TIMING: d_a_cache allocated 4096 MiB on gpu 0. This is the single most important line in the entire log. Even though worker 1 is assigned to GPU 1, the C++ code allocates the d_a_cache—a 4 GB GPU memory buffer—on GPU 0. This is the same GPU-routing bug that was diagnosed for PoRep, now manifesting for SnapDeals.

T+84s: Partition 4 completes synthesis in 84.1 seconds and is sent to the GPU. Worker 0 (assigned to GPU 0) picks it up. Now two partitions are in GPU proving simultaneously: partition 2 on worker 1 (but actually using GPU 0) and partition 4 on worker 0 (also using GPU 0).

T+84s (03:20:17): The crash occurs. A cudaMallocAsync call fails with "out of memory" at supraseal-c2/src/lib.rs:311. The process panics and dumps core.

T+117s (03:20:50): systemd records the process exit with SIGSEGV and reports resource usage: 169.6 GB memory peak, 15.6 GB memory swap peak. The service restarts.

The Critical Clue: d_a_cache allocated 4096 MiB on gpu 0

The line CUZK_TIMING: d_a_cache allocated 4096 MiB on gpu 0 is the Rosetta Stone of this crash. It tells us several things simultaneously:

  1. The multi-GPU bug is still present. Worker 1 is assigned to GPU 1, but the C++ code allocates the d_a_cache on GPU 0. This is the exact same bug that caused PoRep proof corruption on the two-GPU cs-calib host. The C++ groth16_cuda.cu code calls select_gpu(0) or equivalent logic that hard-codes GPU 0 for single-circuit proofs.
  2. Two partitions are competing for GPU 0's VRAM. With the shared mutex hack in place, only one worker can enter the GPU code at a time for partitioned proofs. But the log shows both worker 0 and worker 1 entering GPU proving simultaneously. This means the shared mutex is not working as intended for SnapDeals, or the SnapDeals pipeline uses a different code path that bypasses the mutex. The log shows worker 1 picking up partition 2 and immediately allocating d_a_cache on GPU 0, and worker 0 picking up partition 4 and also entering GPU code. Two workers are simultaneously allocating GPU memory on the same device.
  3. The VRAM budget is exhausted. The RTX 4000 Ada has 20 GB of VRAM. The d_a_cache alone consumes 4 GB. The PCE circuit representation is 15.8 GB (though this is host memory, not GPU memory). The GPU proving kernels require additional memory for the proof itself, intermediate buffers, and the CUDA context. With two concurrent allocations on the same GPU, the 20 GB budget is quickly exceeded. The user's summary says "Snap seems to OOM vram immediately post-synth," which is accurate but incomplete. The OOM is not an inherent property of SnapDeals proving—it is a consequence of two partitions being routed to the same GPU simultaneously, each trying to allocate overlapping memory regions. On a single-GPU system or with proper GPU routing, only one partition would be proving at a time, and the VRAM might be sufficient.

Why WindowPoSt Works but SnapDeals Does Not

The user notes that WindowPoSt works fine on the same hardware. This is an important data point that helps narrow the root cause. WindowPoSt and SnapDeals are different proof types with different circuit sizes and different proving characteristics. The WindowPoSt circuit is smaller, with fewer constraints and a smaller PCE representation. It fits comfortably within the 20 GB VRAM budget even with the GPU-routing bug.

SnapDeals, by contrast, uses the snap-32g circuit, which the PCE extraction reveals to have over 81 million constraints and a 15.8 GB memory footprint. This is a much larger circuit that pushes against the limits of the 20 GB GPU. When the GPU-routing bug causes two partitions to compete for the same GPU's memory, the OOM is inevitable.

The user's framing—"WindowPoSt works fine but Snap seems to OOM"—is a reasonable first-order diagnosis, but it risks misattributing the cause. The OOM is not a SnapDeals-specific problem; it is a multi-GPU routing problem that happens to manifest first on SnapDeals because SnapDeals has larger memory requirements. If the GPU-routing bug were fixed so that each worker uses its assigned GPU, the memory pressure would be halved (each GPU handling one partition instead of both partitions on one GPU), and the SnapDeals workload might succeed.

Assumptions Embedded in the Message

The user's message carries several implicit assumptions:

  1. The crash is a VRAM capacity issue. The user's summary explicitly states "OOM vram immediately post-synth," framing the problem as insufficient GPU memory for SnapDeals. This is a reasonable surface-level diagnosis, but it does not account for the fact that two partitions are simultaneously allocating memory on the same GPU.
  2. SnapDeals and WindowPoSt are comparable workloads. The user contrasts the two to highlight that SnapDeals is uniquely problematic. While this contrast is useful for debugging, it obscures the fact that the underlying GPU-routing bug affects both proof types equally—WindowPoSt just happens to fit within the available memory even when two partitions collide.
  3. The crash is a new issue. The user labels it "New issue," suggesting it is distinct from the PoRep data race being investigated. In reality, the SnapDeals OOM and the PoRep proof corruption share a common root cause: the C++ code's failure to respect the GPU assignment from the Rust engine.
  4. The log contains sufficient diagnostic information. The user provides a raw log dump without additional commentary, assuming the assistant can extract the relevant signal. This is a reasonable assumption for an AI assistant with systems programming knowledge, but it places the burden of interpretation entirely on the receiver.

Input Knowledge Required to Understand This Message

To fully parse this message, a reader needs substantial domain knowledge spanning multiple layers of the system:

Filecoin proving architecture: Understanding what SnapDeals and WindowPoSt are, how they differ, and why they have different circuit sizes. SnapDeals is a proof type used for sector-update operations in the Filecoin network, while WindowPoSt (Proof of Spacetime) is used for ongoing sector maintenance.

CuZK proving engine: Knowledge of the CuZK daemon architecture, including the pipeline stages (synthesis, GPU proving, proof assembly), the role of PCE (Pre-Compiled Circuit Encoding) for accelerating synthesis, and the worker model where multiple GPU workers are assigned to specific GPU devices.

GPU memory management: Understanding of CUDA memory allocation, the d_a_cache mechanism, and how VRAM is consumed during Groth16 proving. The 4 GB d_a_cache allocation is a significant fraction of the 20 GB budget.

The multi-GPU bug context: The preceding conversation establishes that the C++ GPU code always routes single-circuit proofs to GPU 0. Without this context, the d_a_cache allocated 4096 MiB on gpu 0 line for a worker assigned to GPU 1 would appear to be a harmless logging inconsistency rather than a critical bug.

Constraint system representation: Understanding what the PCE extraction statistics mean—81 million constraints, 443 million non-zero entries, 15.8 GB memory footprint—and how these numbers translate to GPU memory requirements during proving.

Output Knowledge Created by This Message

Despite its brevity, this message creates substantial diagnostic knowledge:

  1. SnapDeals PCE characteristics are now known. The extraction log provides precise statistics for the snap-32g circuit: 5 inputs, 80,978,488 aux variables, 81,049,504 constraints, and a 15.8 GB memory footprint. This data is essential for capacity planning and for understanding why the circuit stresses the 20 GB GPU.
  2. The multi-GPU bug affects SnapDeals, not just PoRep. The d_a_cache allocation on GPU 0 for a worker assigned to GPU 1 proves that the GPU-routing bug is not specific to PoRep—it affects all single-circuit proof types. This generalizes the bug and raises its priority.
  3. The shared mutex hack is insufficient for SnapDeals. The log shows two workers entering GPU proving simultaneously despite the shared mutex. This indicates either that the SnapDeals pipeline uses a different code path that bypasses the mutex, or that the mutex is not being acquired correctly for SnapDeals partitioned proofs. Either way, the shared mutex approach is not a viable long-term solution.
  4. The crash provides a baseline for measuring the fix. Once the gpu_index threading fix is deployed, the same SnapDeals workload should succeed because each worker will use its assigned GPU, eliminating the double-allocation on GPU 0. The crash log serves as a "before" measurement.
  5. System resource limits are documented. The systemd summary shows a 169.6 GB host memory peak and 15.6 GB swap peak, indicating that the proving process is also memory-intensive on the CPU side. This is relevant for overall system sizing.

The Thinking Process: What the Assistant Would Infer

An experienced systems engineer reading this log would follow a specific reasoning chain:

First, they would note the PCE extraction statistics and recognize that the SnapDeals circuit is unusually large. The 15.8 GB memory footprint means that even a single partition's proving data is substantial.

Second, they would spot the d_a_cache allocated 4096 MiB on gpu 0 line and immediately recognize the multi-GPU routing bug. This line is the direct link between the SnapDeals crash and the PoRep data race that was already diagnosed. The engineer would think: "This is the same bug. Worker 1 is supposed to use GPU 1, but the C++ code is allocating on GPU 0."

Third, they would correlate the two-worker concurrency with the OOM. Two partitions, each requiring significant GPU memory, are both allocating on GPU 0. The 20 GB VRAM budget is insufficient for two concurrent proving operations on the same device. The engineer would calculate: 4 GB d_a_cache per partition × 2 partitions = 8 GB just for the cache, plus the proving kernels' working sets, plus the CUDA context overhead. The total likely exceeds 20 GB.

Fourth, they would reconcile the user's observation that WindowPoSt works. The WindowPoSt circuit is smaller, so even with two partitions colliding on GPU 0, the total memory stays within budget. This confirms that the root cause is the GPU-routing bug, not an inherent SnapDeals memory requirement.

Fifth, they would recognize that the shared mutex hack is not preventing concurrency for SnapDeals. This might be because SnapDeals uses a different pipeline path (the logs show synth_single{proof_kind=snap-update} and synthesize_snap_deals_partition), or because the mutex was only applied to PoRep paths. Either way, the hack is insufficient.

Finally, they would conclude that the gpu_index threading fix—already in progress—is the correct solution for both the PoRep data race and the SnapDeals OOM. The fix must be deployed to p-dev-ngw-1 as well as cs-calib.

Mistakes and Incorrect Assumptions

The user's message, while valuable, contains some potential misdiagnoses:

"Snap seems to OOM vram immediately post-synth" is a correct observation of the symptom but an incomplete explanation of the cause. The OOM is not a simple capacity exhaustion—it is a consequence of two partitions being routed to the same GPU. If the GPU routing were correct, each GPU would handle one partition, and the memory pressure would be distributed.

The implicit assumption that this is a separate issue from the PoRep bug is incorrect. Both failures stem from the same root cause: the C++ code's hard-coded GPU 0 routing. The PoRep bug manifested as proof corruption (because two workers wrote to the same GPU memory simultaneously), while the SnapDeals bug manifested as OOM (because two workers allocated overlapping memory regions on the same GPU). Different symptoms, same disease.

The assumption that WindowPoSt working implies SnapDeals should also work is a reasonable baseline but ignores the significant difference in circuit size. The 15.8 GB PCE footprint for SnapDeals versus the presumably smaller WindowPoSt circuit means that SnapDeals has much less margin for error in memory management.

Conclusion: A Message That Changed the Trajectory

Message 430 is a pivotal moment in the conversation. It arrives just as the assistant is implementing a multi-file, multi-language fix for a bug that was thought to be PoRep-specific. The SnapDeals crash log proves that the bug is general: it affects all single-circuit proof types, and its symptoms vary depending on the circuit size and the available VRAM.

The message transforms the debugging effort from a narrow investigation of PoRep proof corruption into a broader architectural correction. The gpu_index threading fix, which was being implemented for PoRep, is now understood to be essential for SnapDeals as well. The shared mutex hack, which was already recognized as suboptimal, is now revealed to be insufficient even for its intended purpose.

For the reader of this article, message 430 demonstrates the value of raw diagnostic data in systems debugging. A single line—d_a_cache allocated 4096 MiB on gpu 0—provides the critical link between two seemingly unrelated failures. It shows that the same architectural flaw can manifest as data corruption on one system and out-of-memory on another, depending on the workload and hardware configuration. The message is a reminder that in complex distributed systems, the symptom is not always the disease, and the most valuable diagnostic data is often the data that connects one failure to another.