Skip to content

Graph capture fix - #17

Merged
Ra4ster merged 12 commits into
mainfrom
GraphCaptureFix
Oct 4, 2026
Merged

Ra4ster merged 12 commits into
mainfrom
GraphCaptureFix

Conversation

@Ra4ster

@Ra4ster Ra4ster commented Oct 4, 2026

Copy link
Copy Markdown
Collaborator

Description

Closes CUTLASS integration.

Type of change

  • Bug fix
  • New feature
  • Performance improvement
  • Documentation
  • Build/tooling
  • Breaking change (existing C++ or Python API behavior changes)

Testing

  • python build.py Release builds and DeepityTests passes
  • python build.py Debug builds and DeepityTests passes (required for anything correctness-sensitive)
  • pyright passes (if you touched bindings/, pydeepity/, examples/, or experiments/)

Platform(s) and compiler(s) tested:

  • RHEL Server, Compliant GCC/G++ Compiler
  • AMD EPYC 7643 48-Core
  • NVIDIA A100 PCIE 40GB

Performance impact

Appears to yield <1s speedup. Testing not done, but allocation is safe.

Checklist

  • I've read CONTRIBUTING.md, including the portability rules and numerical gotchas relevant to this change.
  • I've added a CHANGELOG.md entry under [Unreleased] (for any user-facing change).
  • I've updated relevant docs (README, Doxygen comments) if behavior or the public API changed.
  • New/changed code follows the project's .clang-format style.

Notes for reviewers

This is the test run result:

(.venv) [rubaolee2008@a0336 Deepity]$ python mnist.py
Fetching canonical MNIST dataset (idx-ubyte, matching ngc-learn exactly)...

Building network (784->512->10), seed=7...

*** FULL DKP-PC RUN ***
Training DKPPCN: 50 epochs, inference_steps=2, 
lr=0.00373, ir=0.15, lmbda=0.0001, decay_rate=0.94...

Epoch 1/50 | Time: 0.6s | Acc: 90.56% | Avg energy: 167.4288
Epoch 2/50 | Time: 0.8s | Acc: 94.24% | Avg energy: 30.0815
Epoch 3/50 | Time: 1.0s | Acc: 94.88% | Avg energy: 25.1540
Epoch 4/50 | Time: 1.2s | Acc: 94.96% | Avg energy: 21.1647
Epoch 5/50 | Time: 1.4s | Acc: 95.72% | Avg energy: 19.4039
Epoch 6/50 | Time: 1.6s | Acc: 96.92% | Avg energy: 18.6647
Epoch 7/50 | Time: 1.8s | Acc: 96.96% | Avg energy: 17.6914
Epoch 8/50 | Time: 1.9s | Acc: 96.76% | Avg energy: 16.8292
Epoch 9/50 | Time: 2.1s | Acc: 97.24% | Avg energy: 16.2623
Epoch 10/50 | Time: 2.3s | Acc: 97.80% | Avg energy: 15.2082
Epoch 11/50 | Time: 2.5s | Acc: 97.72% | Avg energy: 14.6903
Epoch 12/50 | Time: 2.7s | Acc: 97.56% | Avg energy: 15.2717
Epoch 13/50 | Time: 2.9s | Acc: 97.96% | Avg energy: 14.4918
Epoch 14/50 | Time: 3.1s | Acc: 98.08% | Avg energy: 13.4874
Epoch 15/50 | Time: 3.2s | Acc: 98.00% | Avg energy: 13.2626
Epoch 16/50 | Time: 3.4s | Acc: 97.96% | Avg energy: 13.0047
Epoch 17/50 | Time: 3.6s | Acc: 98.36% | Avg energy: 12.8384
Epoch 18/50 | Time: 3.8s | Acc: 98.48% | Avg energy: 12.1972
Epoch 19/50 | Time: 4.0s | Acc: 98.00% | Avg energy: 12.0181
Epoch 20/50 | Time: 4.2s | Acc: 98.80% | Avg energy: 11.1621
Epoch 21/50 | Time: 4.3s | Acc: 98.96% | Avg energy: 10.9050
Epoch 22/50 | Time: 4.5s | Acc: 98.28% | Avg energy: 10.2527
Epoch 23/50 | Time: 4.7s | Acc: 98.64% | Avg energy: 10.2167
Epoch 24/50 | Time: 4.9s | Acc: 98.52% | Avg energy: 9.6508
Epoch 25/50 | Time: 5.1s | Acc: 98.52% | Avg energy: 9.8337
Epoch 26/50 | Time: 5.3s | Acc: 98.32% | Avg energy: 9.1208
Epoch 27/50 | Time: 5.5s | Acc: 98.76% | Avg energy: 9.1744
Epoch 28/50 | Time: 5.6s | Acc: 99.16% | Avg energy: 8.5406
Epoch 29/50 | Time: 5.8s | Acc: 98.88% | Avg energy: 8.1217
Epoch 30/50 | Time: 6.0s | Acc: 99.04% | Avg energy: 8.0488
Epoch 31/50 | Time: 6.2s | Acc: 98.88% | Avg energy: 7.7218
Epoch 32/50 | Time: 6.4s | Acc: 99.24% | Avg energy: 7.4119
Epoch 33/50 | Time: 6.6s | Acc: 99.24% | Avg energy: 7.3986
Epoch 34/50 | Time: 6.8s | Acc: 99.24% | Avg energy: 6.9620
Epoch 35/50 | Time: 6.9s | Acc: 99.00% | Avg energy: 6.6912
Epoch 36/50 | Time: 7.1s | Acc: 99.36% | Avg energy: 6.6574
Epoch 37/50 | Time: 7.3s | Acc: 99.36% | Avg energy: 6.4908
Epoch 38/50 | Time: 7.5s | Acc: 99.24% | Avg energy: 6.1422
Epoch 39/50 | Time: 7.7s | Acc: 99.44% | Avg energy: 5.9899
Epoch 40/50 | Time: 7.9s | Acc: 99.44% | Avg energy: 5.9080
Epoch 41/50 | Time: 8.1s | Acc: 98.80% | Avg energy: 5.7114
Epoch 42/50 | Time: 8.2s | Acc: 99.40% | Avg energy: 5.5523
Epoch 43/50 | Time: 8.4s | Acc: 98.80% | Avg energy: 5.3992
Epoch 44/50 | Time: 8.6s | Acc: 99.24% | Avg energy: 5.3381
Epoch 45/50 | Time: 8.8s | Acc: 99.40% | Avg energy: 5.2862
Epoch 46/50 | Time: 9.0s | Acc: 99.28% | Avg energy: 5.1409
Epoch 47/50 | Time: 9.2s | Acc: 99.12% | Avg energy: 5.0743
Epoch 48/50 | Time: 9.4s | Acc: 99.08% | Avg energy: 5.0316
Epoch 49/50 | Time: 9.5s | Acc: 99.44% | Avg energy: 4.9311
Epoch 50/50 | Time: 9.7s | Acc: 99.40% | Avg energy: 4.8432

Training complete in 9.7s.

Running final test evaluation...

=== Result ===
DKPPCN test accuracy: 97.81%
Train time: 9.7s

Ra4ster and others added 12 commits October 2, 2026 19:46
…memory.

Graph capture on ConvPCNetwork/SimpleConvPCNetwork/GaussSeidelPCNetwork/
DiscriminativePCNetwork surfaced two pre-existing CUDA backend bugs that
had never been reachable before (nothing used capture on this path):

1. ComputePrecisionWeightedErrorAndEnergy/ComputeErrorAndEnergy both do a
   blocking cudaStreamSynchronize() to read energy back to a host float,
   illegal while a stream is mid-capture. FullPCLayer/SimplePCLayer
   already had a needEnergy bool to skip this; ConvPCLayer/
   SimpleConvPCLayer/GaussSeidelPCLayer/DiscriminativePCLayer never did.

2. ConvPCLayer/DiscriminativePCLayer's input-layer branch read log_p[i]
   directly from host code -- illegal once log_p is a GPU device pointer
   (DeviceMemoryArena), independent of capture entirely.

Added the missing ComputePrecisionWeightedError (no-energy, no-sync) and
Sum (host-safe reduction) backend primitives to CPUBackend/CUDABackend,
threaded needEnergy through all 4 affected layers/networks mirroring
SimplePCLayer's existing pattern, and fixed two nanobind bindings that
became ambiguous once CalculateState had two overloads.

Verified: full CPU rebuild clean, ctest passing, and all 4 TrainStep/
TrainStepWithProjection sanity tests produce bit-for-bit identical energy
trajectories to before this fix -- zero CPU behavioral change, GPU-path
correctness only. The new CUDA kernels themselves are unverified by a
compiler (no GPU on this machine); they closely mirror the existing,
working kernels in the same file.
…e fix).

Reported symptom: on real GPU hardware, the first train_step() call after
switching inference_steps (triggering a graph re-capture) returned the
EXACT same energy as the previous call, as if nothing executed; the next
call then showed roughly double the expected delta, as if the "missing"
update got applied together with the following one. No CUDA error was
printed anywhere in the capture/instantiate/launch sequence.

cudaGraphLaunch() only enqueues the replay -- it returns once the launch
is submitted, not once it completes. Same-stream ordering should already
guarantee the subsequent energy readout waits for it, but added an
explicit IComputeBackend::Synchronize() (CPUBackend: no-op; CUDABackend:
cudaStreamSynchronize()) right after every ReplayGraph() call across all
12 call sites (7 networks) to close any doubt, since this costs nothing
extra on the path that was already about to do its own blocking sync for
the energy value anyway.

Verified: CPU rebuild clean, ctest passing, zero behavioral change on
CPU (Synchronize() is a no-op there). GPU-side effect unverified -- no
GPU on this machine -- this is a hypothesis-driven fix pending retest.
GPU branch, to directly observe graphCaptured/capturedInferenceSteps
state transitions -- diagnosing a reproducible bug where the first
train_step() call after switching inference_steps (triggering recapture)
produces an unchanged energy value, and the NEXT call's delta is
consistently ~2x a full old-step-count delta, not the new one (observed
across two independent runs with different random weights). Sanitizer
(synccheck) reports 0 errors, so this isn't a CUDA API misuse bug -- the
recapture logic itself needs to be observed directly. Temporary, to be
removed once the root cause is found.
…) call

Distinguishes two remaining hypotheses for the GraphCaptureFix bug: whether
the first replay after a fresh cudaGraphInstantiate() genuinely performs no
GPU work, or whether the GPU work happens but the energy readout reported
by the existing [DIAG] lines is stale/host-cached. Reads layers.front()'s
first weight directly via CopyToHost, bypassing CalculateState() entirely,
before and after the recapture branch's ReplayGraph()+Synchronize(). If
W[0] changes across that call, the GPU executed real work and the bug is
in how energy is read out; if it doesn't, the first launch truly is a
no-op and the bug is in graph execution itself.

Temporary, to be removed once the root cause is found.
…resh energy read

The W[0] diagnostic from the last run proved the GPU is NOT idle on the
first replay after a fresh cudaGraphInstantiate(): weights move by a
normal, in-trend amount on every single call, capture or not. So the bug
isn't graph execution -- it's specifically that the energy value returned
by TrainStep() is stale (bit-identical to the previous call's) on the one
call that just did a (re)capture.

Root cause: ConvPCLayer::ComputeMuOnly() caches mu for clamped layers
(isClamped && muCacheValid) to avoid recomputing an input that hasn't
changed. ClampState() normally resets muCacheValid=false at the top of
every TrainStep() call, which is what forces the final, uncaptured
CalculateState() read (the one that produces the returned/printed energy)
to recompute mu from this call's POST-weight-update state. But a (re)capture
runs the whole settling loop as real C++ while *recording* the graph --
that loop's last ComputeMuOnly() call leaves muCacheValid=true as a side
effect, and nothing clears it again before ReplayGraph() runs and the
fresh energy read happens. That read then wrongly takes the cached
branch, returning mu as it was captured (from this call's PRE-update
weights) instead of the fresh value ReplayGraph() just produced.

Fix: explicitly invalidate every layer's mu-cache right after a successful
EndGraphCapture(), mirroring what a fresh ClampState() call would have
done. Diagnostic instrumentation (std::cerr) left in place for this round
of testing; to be removed once confirmed.
… is fixed and confirmed

temp2.py's recapture test now reports RESULT: OK on the GPU node. The
InvalidateMuCache() fix from the previous commit stands on its own.
…capturing network

ConvPCNetwork's confirmed fix (commit 663262a) generalizes: every network
that pairs CUDA graph capture with a clamped-layer mu-cache
(isClamped && muCacheValid, used to skip recomputing an input that hasn't
changed) has the identical latent bug. A (re)capture runs its settling
loop as real C++ while recording the graph, which leaves muCacheValid=true
behind on whichever clamped layer it last touched; nothing clears that
before ReplayGraph() runs and the following fresh, uncaptured energy read
happens, so that read wrongly reuses mu from before this call's own
weight update instead of the value the replay just produced.

Covered now, all capture sites (including TrainStepWithProjection()
variants where present):
  - ConvPCNetwork (TrainStepWithProjection(), missed in the previous commit)
  - SimplePCNetwork
  - SimpleConvPCNetwork (both TrainStep() and TrainStepWithProjection())
  - DiscriminativePCNetwork (both TrainStep() and TrainStepWithProjection())
  - DirectKPPCNetwork
  - FullPCNetwork, FullConvPCNetwork: InvalidateMuCache() already existed
    here (added earlier for an unrelated reason -- iPC's weights changing
    mid-settle), but was only called when useIPC is on; added the same
    unconditional call made everywhere else so the non-iPC path is covered
    too.

GaussSeidelPCNetwork is NOT affected: GaussSeidelPCLayer has no mu-cache
at all, so there's nothing for a recapture to leave stale.

Added InvalidateMuCache() to SimplePCLayer, SimpleConvPCLayer,
DiscriminativePCLayer, and DirectKPPCLayer (ConvPCLayer already has it
from the previous commit; FullPCLayer/FullConvPCLayer already had it for
the iPC reason above).

Builds clean and passes ctest locally (CPU-only, no GPU on this machine);
the GPU-side fix still needs verifying on the cluster the same way
ConvPCNetwork's was, since none of this has been exercised against real
graph capture yet.
@Ra4ster Ra4ster self-assigned this Oct 4, 2026
@Ra4ster Ra4ster added the enhancement New feature or request label Oct 4, 2026
@Ra4ster Ra4ster linked an issue Oct 4, 2026 that may be closed by this pull request
2 tasks done
@Ra4ster
Ra4ster merged commit b5276b1 into main Oct 4, 2026
12 checks passed
@Ra4ster
Ra4ster deleted the GraphCaptureFix branch October 4, 2026 17:11
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

enhancement New feature or request

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[Feature]: Greater CUTLASS Integration

1 participant