From 50482e0dee0c376a8a3f638d807ffa1854ea7ae3 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Fri, 2 Oct 2026 19:46:08 -0400 Subject: [PATCH 01/12] Initial commit on graph fix. --- bindings/NetworkBindings.cpp | 38 +++++- include/deepity/networks/ConvPCNetwork.h | 9 ++ .../networks/DiscriminativePCNetwork.h | 8 ++ .../deepity/networks/GaussSeidelPCNetwork.h | 8 ++ .../deepity/networks/SimpleConvPCNetwork.h | 8 ++ src/networks/ConvPCNetwork.cpp | 125 ++++++++++++++++-- src/networks/DiscriminativePCNetwork.cpp | 117 ++++++++++++++-- src/networks/FullConvPCNetwork.cpp | 58 +++++++- src/networks/GaussSeidelPCNetwork.cpp | 122 +++++++++++++++-- src/networks/SimpleConvPCNetwork.cpp | 113 ++++++++++++++-- temp.py | 14 ++ ...DiscriminativePCNetworkTrainStepSanity.cpp | 75 +++++++++++ .../tGaussSeidelPCNetworkTrainStepSanity.cpp | 83 ++++++++++++ tests/tMatMulLargeAsymmetricCPUVerify.cpp | 57 ++++++++ tests/tSimpleConvPCNetworkTrainStepSanity.cpp | 78 +++++++++++ 15 files changed, 858 insertions(+), 55 deletions(-) create mode 100644 tests/tDiscriminativePCNetworkTrainStepSanity.cpp create mode 100644 tests/tGaussSeidelPCNetworkTrainStepSanity.cpp create mode 100644 tests/tMatMulLargeAsymmetricCPUVerify.cpp create mode 100644 tests/tSimpleConvPCNetworkTrainStepSanity.cpp diff --git a/bindings/NetworkBindings.cpp b/bindings/NetworkBindings.cpp index ca510cb..a57610c 100644 --- a/bindings/NetworkBindings.cpp +++ b/bindings/NetworkBindings.cpp @@ -104,7 +104,18 @@ void bind_networks(nb::module_& m) auto discNetCls = nb::class_( m, "DiscriminativePCNetwork", "Predictive Coding Network."); BindCommonPCNetwork(discNetCls, "DiscriminativePCNetwork"); - discNetCls.def(nb::init<>(), "Construct a network with automatic batch-size detection.") + discNetCls + .def( + "__init__", + [](Deep::DiscriminativePCNetwork* self, const std::string& device) + { + Deep::DeviceType dt = (device == "cuda" || device == "gpu") + ? Deep::DeviceType::DEVICE_GPU + : Deep::DeviceType::DEVICE_CPU; + new (self) Deep::DiscriminativePCNetwork(dt); + }, + nb::arg("device") = "cpu", + "Construct a network with automatic batch-size detection, on the given device.") .def( "add_layer", [](Deep::DiscriminativePCNetwork& self, @@ -298,7 +309,17 @@ void bind_networks(nb::module_& m) nb::class_( m, "GaussSeidelPCNetwork", "Predictive Coding Network with Gauss-Seidel settling dynamics.") - .def(nb::init(), nb::arg("batch_size")) + .def( + "__init__", + [](Deep::GaussSeidelPCNetwork* self, int batch_size, const std::string& device) + { + Deep::DeviceType dt = (device == "cuda" || device == "gpu") + ? Deep::DeviceType::DEVICE_GPU + : Deep::DeviceType::DEVICE_CPU; + new (self) Deep::GaussSeidelPCNetwork(batch_size, dt); + }, + nb::arg("batch_size"), + nb::arg("device") = "cpu") .def( "add_layer", [](Deep::GaussSeidelPCNetwork& self, @@ -947,7 +968,18 @@ void bind_networks(nb::module_& m) }); nb::class_(m, "ConvPCNetwork", "Convolutional Predictive Coding Network.") - .def(nb::init(), nb::arg("batch_size"), "Construct a network with a fixed batch size.") + .def( + "__init__", + [](Deep::ConvPCNetwork* self, int batch_size, const std::string& device) + { + Deep::DeviceType dt = (device == "cuda" || device == "gpu") + ? Deep::DeviceType::DEVICE_GPU + : Deep::DeviceType::DEVICE_CPU; + new (self) Deep::ConvPCNetwork(batch_size, dt); + }, + nb::arg("batch_size"), + nb::arg("device") = "cpu", + "Construct a network with a fixed batch size, on the given device.") .def( "add_layer", [](Deep::ConvPCNetwork& self, diff --git a/include/deepity/networks/ConvPCNetwork.h b/include/deepity/networks/ConvPCNetwork.h index 49bbc57..0b83169 100644 --- a/include/deepity/networks/ConvPCNetwork.h +++ b/include/deepity/networks/ConvPCNetwork.h @@ -107,5 +107,14 @@ namespace Deep #endif int batchSize; OptimizerType pendingOpt = OptimizerType::SGD; + + // CUDA graph capture state, GPU-only: TrainStep() and + // TrainStepWithProjection() capture different op sequences, so + // each needs its own capture-validity tracking (a graph captured + // for one is not valid for the other). + bool graphCaptured = false; + int capturedInferenceSteps = -1; + bool graphCapturedWithProjection = false; + int capturedInferenceStepsWithProjection = -1; }; } // namespace Deep diff --git a/include/deepity/networks/DiscriminativePCNetwork.h b/include/deepity/networks/DiscriminativePCNetwork.h index 550738a..67ca8a3 100644 --- a/include/deepity/networks/DiscriminativePCNetwork.h +++ b/include/deepity/networks/DiscriminativePCNetwork.h @@ -233,5 +233,13 @@ namespace Deep std::unique_ptr gpuArena; #endif friend class PCNDiagnostics; + + // CUDA graph capture state, GPU-only: TrainStep() and + // TrainStepWithProjection() capture different op sequences, so + // each needs its own capture-validity tracking. + bool graphCaptured = false; + int capturedInferenceSteps = -1; + bool graphCapturedWithProjection = false; + int capturedInferenceStepsWithProjection = -1; }; } diff --git a/include/deepity/networks/GaussSeidelPCNetwork.h b/include/deepity/networks/GaussSeidelPCNetwork.h index f45ca37..0b87298 100644 --- a/include/deepity/networks/GaussSeidelPCNetwork.h +++ b/include/deepity/networks/GaussSeidelPCNetwork.h @@ -162,5 +162,13 @@ namespace Deep #endif /// @brief The batch size shared by every layer in the network. int batchSize; + + // CUDA graph capture state, GPU-only: TrainStep() and + // TrainStepWithProjection() capture different op sequences, so + // each needs its own capture-validity tracking. + bool graphCaptured = false; + int capturedInferenceSteps = -1; + bool graphCapturedWithProjection = false; + int capturedInferenceStepsWithProjection = -1; }; } diff --git a/include/deepity/networks/SimpleConvPCNetwork.h b/include/deepity/networks/SimpleConvPCNetwork.h index b3408af..e9e9890 100644 --- a/include/deepity/networks/SimpleConvPCNetwork.h +++ b/include/deepity/networks/SimpleConvPCNetwork.h @@ -118,5 +118,13 @@ namespace Deep int batchSize; OptimizerType pendingOpt = OptimizerType::SGD; friend class PCNDiagnostics; + + // CUDA graph capture state, GPU-only: TrainStep() and + // TrainStepWithProjection() capture different op sequences, so + // each needs its own capture-validity tracking. + bool graphCaptured = false; + int capturedInferenceSteps = -1; + bool graphCapturedWithProjection = false; + int capturedInferenceStepsWithProjection = -1; }; } // namespace Deep \ No newline at end of file diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index 0e7e4d8..c9c5d07 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -1,5 +1,6 @@ #include #include +#include namespace Deep { @@ -116,14 +117,61 @@ namespace Deep Clamp(x); GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) + if (device == DeviceType::DEVICE_GPU) { - finalEnergy = CalculateState(); - UpdateState(); + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCaptured = true; + capturedInferenceSteps = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + // A fresh, uncaptured call: a value returned by a call made + // INSIDE the captured region above would be stale on every + // ReplayGraph() after the first (only the GPU kernel launches + // are replayed, not the host-side C++ call that returned this + // float), so the real energy readout always happens out here, + // identically for both the CPU and GPU paths. + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -172,17 +220,72 @@ namespace Deep { ResetState(); Clamp(x); - ProjectForward(); + // Clamped BEFORE ProjectForward() (reordered from the original + // ProjectForward()-then-clamp), required for two reasons: (1) so + // ProjectForward()'s own IsClamped() guard actually protects the + // terminal layer from being overwritten by a forward-projected + // guess, and (2) so ProjectForward() can safely move inside the + // captured region below -- ClampState()'s CopyFromHost(z, y.data(), + // ...) call captures a SPECIFIC host pointer; replaying that + // inside a captured graph with a different batch's `y` (a + // different vector, different address, every call) would copy + // stale data. Keeping it out here means it genuinely re-executes + // with fresh data on every call, same pattern DirectKPPCNetwork + // already uses. GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) + if (device == DeviceType::DEVICE_GPU) { - finalEnergy = CalculateState(); - UpdateState(); + if (!graphCapturedWithProjection || capturedInferenceStepsWithProjection != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCapturedWithProjection = true; + capturedInferenceStepsWithProjection = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/src/networks/DiscriminativePCNetwork.cpp b/src/networks/DiscriminativePCNetwork.cpp index f449f45..92c913d 100644 --- a/src/networks/DiscriminativePCNetwork.cpp +++ b/src/networks/DiscriminativePCNetwork.cpp @@ -2,6 +2,7 @@ #include #include #include +#include #include #include @@ -123,18 +124,69 @@ namespace Deep ProtectFPU(); ResetState(); Clamp(x); - ProjectForward(); + // Clamped BEFORE ProjectForward() (reordered from the original + // ProjectForward()-then-clamp): required so ProjectForward()'s own + // IsClamped() guard protects the terminal layer, and so + // ProjectForward() can safely move inside the captured region + // below -- see ConvPCNetwork::TrainStepWithProjection()'s identical + // comment for why. GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) + if (device == DeviceType::DEVICE_GPU) { - finalEnergy = CalculateState(); - UpdateState(); + if (!graphCapturedWithProjection || capturedInferenceStepsWithProjection != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdatePrecision(); + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCapturedWithProjection = true; + capturedInferenceStepsWithProjection = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdatePrecision(); + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdatePrecision(); + UpdateWeights(); } - UpdatePrecision(); - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -168,14 +220,55 @@ namespace Deep Clamp(x); GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; t++) + if (device == DeviceType::DEVICE_GPU) { - finalEnergy = CalculateState(); - UpdateState(); + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + for (int t = 0; t < inferenceSteps; t++) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCaptured = true; + capturedInferenceSteps = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + } + else + { + for (int t = 0; t < inferenceSteps; t++) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; t++) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/src/networks/FullConvPCNetwork.cpp b/src/networks/FullConvPCNetwork.cpp index 48adb53..7b3df35 100644 --- a/src/networks/FullConvPCNetwork.cpp +++ b/src/networks/FullConvPCNetwork.cpp @@ -1,6 +1,7 @@ #include #include #include +#include #include #include #include @@ -146,13 +147,56 @@ float FullConvPCNetwork::TrainStep(const std::vector& x, const std::vecto } }; - ProjectForward(); - CalculateTerminalError(); - DirectFeedbackUpdate(); - for (int t = 0; t < inferenceSteps; t++) - settleStep(); - if (!useIPC) - UpdateWeights(); + if (device == DeviceType::DEVICE_GPU) + { + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); + CalculateTerminalError(); + DirectFeedbackUpdate(); + for (int t = 0; t < inferenceSteps; t++) + settleStep(); + if (!useIPC) + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCaptured = true; + capturedInferenceSteps = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + } + else + { + ProjectForward(); + CalculateTerminalError(); + DirectFeedbackUpdate(); + for (int t = 0; t < inferenceSteps; t++) + settleStep(); + if (!useIPC) + UpdateWeights(); + } + } + else + { + ProjectForward(); + CalculateTerminalError(); + DirectFeedbackUpdate(); + for (int t = 0; t < inferenceSteps; t++) + settleStep(); + if (!useIPC) + UpdateWeights(); + } float finalEnergy = 0.0f; for (auto& l : layers) diff --git a/src/networks/GaussSeidelPCNetwork.cpp b/src/networks/GaussSeidelPCNetwork.cpp index 2d3a0da..575c101 100644 --- a/src/networks/GaussSeidelPCNetwork.cpp +++ b/src/networks/GaussSeidelPCNetwork.cpp @@ -1,5 +1,6 @@ #include #include +#include namespace Deep { @@ -100,11 +101,57 @@ namespace Deep Clamp(x); GetTerminalLayer()->ClampState(y); + if (device == DeviceType::DEVICE_GPU) + { + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + for (int t = 0; t < inferenceSteps; ++t) + Step(); + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCaptured = true; + capturedInferenceSteps = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + Step(); + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + Step(); + UpdateWeights(); + } + + // Fresh, uncaptured energy readout: sweeps 2+3 of Step() (recompute + // predictions/errors from the CURRENT z), deliberately without + // sweep 1's UpdateState() -- calling the full Step() again here + // would apply an extra, unintended settling step beyond + // inferenceSteps, on top of the same stale-return-value-under- + // graph-capture concern every other ported network has. float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) - finalEnergy = Step(); + for (auto &l : layers) + l->ComputePrediction(); + for (auto &l : layers) + finalEnergy += l->ComputeError(); - UpdateWeights(); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -114,19 +161,72 @@ namespace Deep { ResetState(); Clamp(x); - ProjectForward(); + // Clamped BEFORE ProjectForward() (reordered from the original + // ProjectForward()-then-clamp): required so ProjectForward()'s own + // IsClamped() guard protects the terminal layer, and so + // ProjectForward() can safely move inside the captured region + // below -- see ConvPCNetwork::TrainStepWithProjection()'s identical + // comment for why. The end result for z3/mu3 (and therefore the + // immediate ComputeError() below) is unchanged either way: + // ComputePrediction() for the second-to-last layer still runs + // unconditionally, only the backend->Copy into an ALREADY-clamped + // terminal's z is what gets skipped. GetTerminalLayer()->ClampState(y); - // Only the terminal layer's error is computed immediately (e3 = - // z3 - mu3, target vs. projected prediction); hidden layers stay - // at zero. - GetTerminalLayer()->ComputeError(); + if (device == DeviceType::DEVICE_GPU) + { + if (!graphCapturedWithProjection || capturedInferenceStepsWithProjection != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); + // Only the terminal layer's error is computed immediately + // (e3 = z3 - mu3, target vs. projected prediction); hidden + // layers stay at zero. + GetTerminalLayer()->ComputeError(); + for (int t = 0; t < inferenceSteps; ++t) + Step(); + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCapturedWithProjection = true; + capturedInferenceStepsWithProjection = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + } + else + { + ProjectForward(); + GetTerminalLayer()->ComputeError(); + for (int t = 0; t < inferenceSteps; ++t) + Step(); + UpdateWeights(); + } + } + else + { + ProjectForward(); + GetTerminalLayer()->ComputeError(); + for (int t = 0; t < inferenceSteps; ++t) + Step(); + UpdateWeights(); + } float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) - finalEnergy = Step(); + for (auto &l : layers) + l->ComputePrediction(); + for (auto &l : layers) + finalEnergy += l->ComputeError(); - UpdateWeights(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/src/networks/SimpleConvPCNetwork.cpp b/src/networks/SimpleConvPCNetwork.cpp index 8788294..841086e 100644 --- a/src/networks/SimpleConvPCNetwork.cpp +++ b/src/networks/SimpleConvPCNetwork.cpp @@ -1,5 +1,6 @@ #include #include +#include namespace Deep @@ -107,14 +108,55 @@ namespace Deep Clamp(x); GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) + if (device == DeviceType::DEVICE_GPU) { - finalEnergy = CalculateState(); - UpdateState(); + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCaptured = true; + capturedInferenceSteps = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -163,17 +205,66 @@ namespace Deep { ResetState(); Clamp(x); - ProjectForward(); + // Clamped BEFORE ProjectForward() (reordered from the original + // ProjectForward()-then-clamp): required so ProjectForward()'s own + // IsClamped() guard protects the terminal layer, and so + // ProjectForward() can safely move inside the captured region + // below -- see ConvPCNetwork::TrainStepWithProjection()'s identical + // comment for why. GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) + if (device == DeviceType::DEVICE_GPU) { - finalEnergy = CalculateState(); - UpdateState(); + if (!graphCapturedWithProjection || capturedInferenceStepsWithProjection != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCapturedWithProjection = true; + capturedInferenceStepsWithProjection = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/temp.py b/temp.py index 438a527..c8931e0 100644 --- a/temp.py +++ b/temp.py @@ -135,6 +135,20 @@ def main() -> None: current_fl = FL * (DECAY_RATE ** epoch) net.set_feedback_rate(current_fl) + # Root cause of the late-epoch instability (energy bottoms out, + # then rises again, accuracy declining): ir was never decayed + # here, so as lr/fl made weight updates more conservative over + # epochs, settling kept taking full-size steps -- a growing + # outer/inner speed mismatch that eventually destabilized the + # energy landscape. Confirmed by isolating it: W itself never + # explodes (tracked per-layer Frobenius norms, grows <30% over 35 + # epochs either way); disabling DKP's Psi feedback entirely + # (fl=0) does NOT fix it either, same pattern, slightly earlier + # onset; decaying ir alongside lr/fl does fix it, cleanly + # monotonic energy for the full run, both with and without DKP. + current_ir = IR * (DECAY_RATE ** epoch) + net.set_inference_rate(current_ir) + indices = rng.permutation(len(X_train)) X_shuf, Y_shuf = X_train[indices], Y_train[indices] diff --git a/tests/tDiscriminativePCNetworkTrainStepSanity.cpp b/tests/tDiscriminativePCNetworkTrainStepSanity.cpp new file mode 100644 index 0000000..e9f98c1 --- /dev/null +++ b/tests/tDiscriminativePCNetworkTrainStepSanity.cpp @@ -0,0 +1,75 @@ +/** + * @file tDiscriminativePCNetworkTrainStepSanity.cpp + * @brief Basic sanity check for DiscriminativePCNetwork::TrainStep() and + * TrainStepWithProjection(), verifying the CUDA-graph-capture + * restructuring of both (CPU path only, no GPU hardware available here) + * preserved correct CPU behavior: energy decreases over repeated calls + * on a fixed input/target pair, no crashes or non-finite values. + */ +#include +#include +#include +#include +#include + +using namespace Deep; + +int main() +{ + DiscriminativePCNetwork net(1); + net.AddLayer(4, 8, 1e-3f, 0.1f, 0.01f, 0.0f, ActivationType::TANH, ActivationType::dTANH); + net.AddLayer(8, 3, 1e-3f, 0.1f, 0.01f, 0.0f, ActivationType::LINEAR, ActivationType::dLINEAR); + net.AddLayer(3, 0, 1e-3f, 0.1f, 0.01f, 0.0f, ActivationType::LINEAR, ActivationType::dLINEAR); + net.Compile(); + + std::mt19937 rng(11); + net.RandomizeWeights(rng); + + std::vector x = {0.3f, -0.6f, 0.1f, 0.9f}; + std::vector y = {0.4f, -0.2f, 0.7f}; + + printf("=== TrainStep() ===\n"); + bool pass = true; + float firstTrainStepEnergy = -1.0f, lastTrainStepEnergy = -1.0f; + for (int i = 0; i < 20; ++i) + { + float e = net.TrainStep(x, y, 5); + printf(" step %d: energy=%.6f\n", i, e); + if (!std::isfinite(e)) + { + printf(" FAIL: non-finite energy\n"); + pass = false; + break; + } + if (i == 0) + firstTrainStepEnergy = e; + lastTrainStepEnergy = e; + } + + printf("=== TrainStepWithProjection() ===\n"); + float firstProjEnergy = -1.0f, lastProjEnergy = -1.0f; + for (int i = 0; i < 20; ++i) + { + float e = net.TrainStepWithProjection(x, y, 5); + printf(" step %d: energy=%.6f\n", i, e); + if (!std::isfinite(e)) + { + printf(" FAIL: non-finite energy\n"); + pass = false; + break; + } + if (i == 0) + firstProjEnergy = e; + lastProjEnergy = e; + } + + printf("TrainStep(): %.6f -> %.6f, TrainStepWithProjection(): %.6f -> %.6f\n", + firstTrainStepEnergy, lastTrainStepEnergy, firstProjEnergy, lastProjEnergy); + + // The real check: energy actually decreased over 20 calls, not an + // arbitrary absolute threshold. + pass = pass && lastTrainStepEnergy < firstTrainStepEnergy && lastProjEnergy < firstProjEnergy; + + printf("%s\n", pass ? "PASS" : "FAIL"); + return pass ? 0 : 1; +} diff --git a/tests/tGaussSeidelPCNetworkTrainStepSanity.cpp b/tests/tGaussSeidelPCNetworkTrainStepSanity.cpp new file mode 100644 index 0000000..99a74ff --- /dev/null +++ b/tests/tGaussSeidelPCNetworkTrainStepSanity.cpp @@ -0,0 +1,83 @@ +/** + * @file tGaussSeidelPCNetworkTrainStepSanity.cpp + * @brief Basic sanity check for GaussSeidelPCNetwork::TrainStep() and + * TrainStepWithProjection(), neither of which any existing test actually + * exercises (tGaussSeidelPCLayerVerify.cpp's own SettleNetwork() helper + * bypasses both, calling ResetState/Clamp/ProjectForward/Step() by hand). + * Written to verify the CUDA-graph-capture restructuring of both methods + * (CPU path only, no GPU hardware available here) didn't change their + * CPU behavior: confirms energy decreases over repeated calls on a fixed + * input/target pair, and that nothing crashes or produces non-finite + * values across many calls (the graph-capture code paths are gated on + * device==DEVICE_GPU and never run here, but the restructuring also + * changed the CPU path's own energy-readout timing, see + * ConvPCNetwork::TrainStep()'s comment for why). + */ +#include +#include +#include +#include +#include + +using namespace Deep; + +int main() +{ + GaussSeidelPCNetwork net(1); + net.AddLayer(4, 8, 1e-3f, 0.1f, 0.0f, ActivationType::TANH, ActivationType::dTANH); + net.AddLayer(8, 3, 1e-3f, 0.1f, 0.0f, ActivationType::LINEAR, ActivationType::dLINEAR); + net.AddLayer(3, 0, 1e-3f, 0.1f, 0.0f, ActivationType::LINEAR, ActivationType::dLINEAR); + net.Compile(); + + std::mt19937 rng(11); + net.RandomizeWeights(rng); + + std::vector x = {0.3f, -0.6f, 0.1f, 0.9f}; + std::vector y = {0.4f, -0.2f, 0.7f}; + + printf("=== TrainStep() ===\n"); + bool pass = true; + float firstTrainStepEnergy = -1.0f, lastTrainStepEnergy = -1.0f; + for (int i = 0; i < 20; ++i) + { + float e = net.TrainStep(x, y, 5); + printf(" step %d: energy=%.6f\n", i, e); + if (!std::isfinite(e)) + { + printf(" FAIL: non-finite energy\n"); + pass = false; + break; + } + if (i == 0) + firstTrainStepEnergy = e; + lastTrainStepEnergy = e; + } + + printf("=== TrainStepWithProjection() ===\n"); + float firstProjEnergy = -1.0f, lastProjEnergy = -1.0f; + for (int i = 0; i < 20; ++i) + { + float e = net.TrainStepWithProjection(x, y, 5); + printf(" step %d: energy=%.6f\n", i, e); + if (!std::isfinite(e)) + { + printf(" FAIL: non-finite energy\n"); + pass = false; + break; + } + if (i == 0) + firstProjEnergy = e; + lastProjEnergy = e; + } + + printf("TrainStep(): %.6f -> %.6f, TrainStepWithProjection(): %.6f -> %.6f\n", + firstTrainStepEnergy, lastTrainStepEnergy, firstProjEnergy, lastProjEnergy); + + // The real check: energy actually decreased over 20 calls, not an + // arbitrary absolute threshold (lr here is conservative, so the exact + // magnitude reached in 20 steps isn't a meaningful bar on its own). + pass = pass && lastTrainStepEnergy < firstTrainStepEnergy && lastProjEnergy < firstProjEnergy; + + printf("%s\n", pass ? "PASS" : "FAIL"); + return pass ? 0 : 1; +} diff --git a/tests/tMatMulLargeAsymmetricCPUVerify.cpp b/tests/tMatMulLargeAsymmetricCPUVerify.cpp new file mode 100644 index 0000000..169d06b --- /dev/null +++ b/tests/tMatMulLargeAsymmetricCPUVerify.cpp @@ -0,0 +1,57 @@ +/** + * @file tMatMulLargeAsymmetricCPUVerify.cpp + * @brief CPU-only isolation of tMatMulLargeAsymmetricVerify.cpp's own + * CPU reference call -- that test is gated behind DEEPITY_CUDA_ENABLED + * (its whole point is a CPU-vs-GPU differential check), so it never + * builds or runs on a CUDA-less machine, even though its CPU-side call + * alone needs no GPU at all. Confirmed segfaulting on a plain CPU-only + * build (reported during OSC GPU testing of unrelated work, then + * reproduced here with no CUDA involved): backend->MatMul() with + * transA=true, transB=true, M=1024, N=12288, K=250 crashes before + * returning. Root cause not yet investigated. + */ +#include +#include +#include +#include + +using namespace Deep; + +int main() +{ + const int M = 1024, N = 12288, K = 250; + + std::mt19937 rng(42); + std::normal_distribution dist(0.0f, 1.0f); + + std::vector hProj(K * M); + for (auto &v : hProj) v = dist(rng); + std::vector hZF(K * N); + for (auto &v : hZF) v = dist(rng); + std::vector hWZero(M * N, 0.0f); + + printf("Allocated host buffers OK.\n"); + + auto backend = CreateBackend(DeviceType::DEVICE_CPU); + Tensor proj(backend.get(), DeviceType::DEVICE_CPU, hProj); + Tensor zF(backend.get(), DeviceType::DEVICE_CPU, hZF); + Tensor W(backend.get(), DeviceType::DEVICE_CPU, hWZero); + + printf("Tensors constructed OK. Calling MatMul...\n"); + fflush(stdout); + + backend->MatMul( + /*transA=*/true, /*transB=*/true, + M, N, K, + 0.5f, proj.Data(), M, + zF.Data(), N, + 1.0f, W.Data(), N); + + printf("MatMul returned OK.\n"); + + std::vector result(M * N); + W.CopyToHost(result.data()); + printf("CopyToHost OK. result[0]=%f\n", result[0]); + + return 0; +} diff --git a/tests/tSimpleConvPCNetworkTrainStepSanity.cpp b/tests/tSimpleConvPCNetworkTrainStepSanity.cpp new file mode 100644 index 0000000..af9b0b2 --- /dev/null +++ b/tests/tSimpleConvPCNetworkTrainStepSanity.cpp @@ -0,0 +1,78 @@ +/** + * @file tSimpleConvPCNetworkTrainStepSanity.cpp + * @brief Basic sanity check for SimpleConvPCNetwork::TrainStep() and + * TrainStepWithProjection(), verifying the CUDA-graph-capture + * restructuring of both (CPU path only, no GPU hardware available here) + * preserved correct CPU behavior: energy decreases over repeated calls + * on a fixed input/target pair, no crashes or non-finite values. + */ +#include +#include +#include +#include +#include + +using namespace Deep; + +int main() +{ + SimpleConvPCNetwork net(1); + // 1x4x4 -> conv 3x3 stride1 -> 2 channels, 2x2 -> flatten to terminal. + net.AddLayer(1, 2, 4, 4, 3, 3, 1, 1, 0, 0, 1e-3f, 0.1f, 0.0f, ActivationType::TANH, + ActivationType::dTANH); + net.AddLayer(2, 0, 2, 2, 1, 1, 1, 1, 0, 0, 1e-3f, 0.1f, 0.0f, ActivationType::LINEAR, + ActivationType::dLINEAR); + net.Compile(); + + std::mt19937 rng(11); + net.RandomizeWeights(rng); + + std::vector x = {0.2f, -0.5f, 0.8f, 0.1f, -0.3f, 0.6f, -0.7f, 0.4f, + 0.9f, -0.2f, 0.3f, -0.6f, 0.5f, -0.9f, 0.1f, -0.1f}; + std::vector y = {0.4f, -0.2f, 0.1f, 0.3f, -0.5f, 0.2f, 0.1f, -0.3f}; + + printf("=== TrainStep() ===\n"); + bool pass = true; + float firstTrainStepEnergy = -1.0f, lastTrainStepEnergy = -1.0f; + for (int i = 0; i < 20; ++i) + { + float e = net.TrainStep(x, y, 5); + printf(" step %d: energy=%.6f\n", i, e); + if (!std::isfinite(e)) + { + printf(" FAIL: non-finite energy\n"); + pass = false; + break; + } + if (i == 0) + firstTrainStepEnergy = e; + lastTrainStepEnergy = e; + } + + printf("=== TrainStepWithProjection() ===\n"); + float firstProjEnergy = -1.0f, lastProjEnergy = -1.0f; + for (int i = 0; i < 20; ++i) + { + float e = net.TrainStepWithProjection(x, y, 5); + printf(" step %d: energy=%.6f\n", i, e); + if (!std::isfinite(e)) + { + printf(" FAIL: non-finite energy\n"); + pass = false; + break; + } + if (i == 0) + firstProjEnergy = e; + lastProjEnergy = e; + } + + printf("TrainStep(): %.6f -> %.6f, TrainStepWithProjection(): %.6f -> %.6f\n", + firstTrainStepEnergy, lastTrainStepEnergy, firstProjEnergy, lastProjEnergy); + + // The real check: energy actually decreased over 20 calls, not an + // arbitrary absolute threshold. + pass = pass && lastTrainStepEnergy < firstTrainStepEnergy && lastProjEnergy < firstProjEnergy; + + printf("%s\n", pass ? "PASS" : "FAIL"); + return pass ? 0 : 1; +} From 5ca1e5cfb89636976218cc9912d7a0371a87ca30 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 00:31:08 -0400 Subject: [PATCH 02/12] Fix CUDA graph-capture segfault: illegal sync + host deref of device 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. --- bindings/LayerBindings.cpp | 8 +- include/deepity/backend/CPUBackend.h | 5 ++ include/deepity/backend/CUDABackend.h | 5 ++ include/deepity/backend/IComputeBackend.h | 19 +++++ include/deepity/layers/ConvPCLayer.h | 9 ++- .../deepity/layers/DiscriminativePCLayer.h | 9 ++- include/deepity/layers/GaussSeidelPCLayer.h | 10 ++- include/deepity/layers/SimpleConvPCLayer.h | 8 +- include/deepity/networks/ConvPCNetwork.h | 7 +- .../networks/DiscriminativePCNetwork.h | 8 +- .../deepity/networks/GaussSeidelPCNetwork.h | 11 ++- .../deepity/networks/SimpleConvPCNetwork.h | 6 +- src/backend/CPUBackend.cpp | 19 +++++ src/backend/cuda/CUDABackendFusedOps.cu | 76 +++++++++++++++++-- src/layers/ConvPCLayer.cpp | 22 ++++-- src/layers/DiscriminativePCLayer.cpp | 21 +++-- src/layers/GaussSeidelPCLayer.cpp | 8 +- src/layers/SimpleConvPCLayer.cpp | 8 +- src/networks/ConvPCNetwork.cpp | 22 +++--- src/networks/DiscriminativePCNetwork.cpp | 18 ++--- src/networks/GaussSeidelPCNetwork.cpp | 35 +++++---- src/networks/SimpleConvPCNetwork.cpp | 22 +++--- 22 files changed, 281 insertions(+), 75 deletions(-) diff --git a/bindings/LayerBindings.cpp b/bindings/LayerBindings.cpp index 493835e..9acf7a1 100644 --- a/bindings/LayerBindings.cpp +++ b/bindings/LayerBindings.cpp @@ -270,7 +270,10 @@ void bind_layers(nb::module_& m) nb::arg("activation_deriv") = "drelu"); gsLayerCls.def("update_state", &Deep::GaussSeidelPCLayer::UpdateState) .def("compute_prediction", &Deep::GaussSeidelPCLayer::ComputePrediction) - .def("compute_error", &Deep::GaussSeidelPCLayer::ComputeError) + .def("compute_error", + static_cast( + &Deep::GaussSeidelPCLayer::ComputeError), + nb::arg("need_energy") = true) .def_prop_ro("mu", [](Deep::GaussSeidelPCLayer& self) { @@ -672,7 +675,8 @@ void bind_layers(nb::module_& m) nb::arg("lmbda") = 1e-2f, nb::arg("activation") = "relu", nb::arg("activation_deriv") = "drelu") - .def("calculate_state", &Deep::ConvPCLayer::CalculateState) + .def("calculate_state", + static_cast(&Deep::ConvPCLayer::CalculateState)) .def("update_state", &Deep::ConvPCLayer::UpdateState) .def("update_weights", &Deep::ConvPCLayer::UpdateWeights) .def("update_precision", &Deep::ConvPCLayer::UpdatePrecision) diff --git a/include/deepity/backend/CPUBackend.h b/include/deepity/backend/CPUBackend.h index 1c4a57c..3ce1364 100644 --- a/include/deepity/backend/CPUBackend.h +++ b/include/deepity/backend/CPUBackend.h @@ -54,6 +54,8 @@ class CPUBackend : public IComputeBackend const float* B, int ldb, float beta, float* C, int ldc) noexcept override; /// @copydoc Deep::IComputeBackend::SumRows void SumRows(float* dst, const float* src, size_t batchSize, size_t width) noexcept override; + /// @copydoc Deep::IComputeBackend::Sum + float Sum(const float* buf, size_t n) noexcept override; /// @copydoc Deep::IComputeBackend::Scale void Scale(float* buf, size_t n, float alpha) noexcept override; @@ -153,6 +155,9 @@ class CPUBackend : public IComputeBackend float ComputePrecisionWeightedErrorAndEnergy(float* e, const float* z, const float* mu, const float* p, size_t batchSize, size_t width) noexcept override; + /// @copydoc Deep::IComputeBackend::ComputePrecisionWeightedError + void ComputePrecisionWeightedError(float* e, const float* z, const float* mu, const float* p, + size_t batchSize, size_t width) noexcept override; /// @copydoc Deep::IComputeBackend::AxpyBroadcastInto void AxpyBroadcastInto(float* y, const float* x, const float* factor, size_t batchSize, size_t width, float alpha) noexcept override; diff --git a/include/deepity/backend/CUDABackend.h b/include/deepity/backend/CUDABackend.h index 30195fc..0c761b2 100644 --- a/include/deepity/backend/CUDABackend.h +++ b/include/deepity/backend/CUDABackend.h @@ -45,6 +45,8 @@ class CUDABackend : public IComputeBackend /// @copydoc Deep::IComputeBackend::SumRows void SumRows(float* dst, const float* src, size_t batchSize, size_t width) noexcept override; + /// @copydoc Deep::IComputeBackend::Sum + float Sum(const float* buf, size_t n) noexcept override; /// @brief Must be called once, before Compile()'s first /// BeginGraphCapture(), for any backend that needs to prepare @@ -146,6 +148,9 @@ class CUDABackend : public IComputeBackend float ComputePrecisionWeightedErrorAndEnergy(float* e, const float* z, const float* mu, const float* p, size_t batchSize, size_t width) noexcept override; + /// @copydoc Deep::IComputeBackend::ComputePrecisionWeightedError + void ComputePrecisionWeightedError(float* e, const float* z, const float* mu, const float* p, + size_t batchSize, size_t width) noexcept override; /// @copydoc Deep::IComputeBackend::AxpyBroadcastInto void AxpyBroadcastInto(float* y, const float* x, const float* factor, size_t batchSize, size_t width, float alpha) noexcept override; diff --git a/include/deepity/backend/IComputeBackend.h b/include/deepity/backend/IComputeBackend.h index 2d4c44c..ba3ca43 100644 --- a/include/deepity/backend/IComputeBackend.h +++ b/include/deepity/backend/IComputeBackend.h @@ -103,6 +103,14 @@ class IComputeBackend /// [width,batchSize] with no data movement (verified numerically). virtual void SumRows(float* dst, const float* src, size_t batchSize, size_t width) noexcept = 0; + /// @brief Returns the plain (signed) sum of every element in buf. + /// Synchronous on CUDABackend (reads the reduction back to a host + /// float before returning) -- same constraint as ComputeErrorAndEnergy + /// and friends: never call this from inside a captured CUDA graph + /// region, only from code that runs outside any + /// BeginGraphCapture()/EndGraphCapture() pair. + virtual float Sum(const float* buf, size_t n) noexcept = 0; + // Elementwise scalar ops /// @brief buf[i] *= alpha for all i in [0, n). @@ -337,6 +345,17 @@ class IComputeBackend virtual float ComputePrecisionWeightedErrorAndEnergy(float* e, const float* z, const float* mu, const float* p, size_t batchSize, size_t width) noexcept = 0; + /// @brief Same as ComputePrecisionWeightedErrorAndEnergy(), without + /// computing (or returning) the energy, for callers that only need e. + /// On CUDABackend this additionally means: no blocking + /// cudaStreamSynchronize() -- the ...AndEnergy() variant must + /// synchronously read the energy reduction back to a host float, which + /// is illegal while a CUDA graph is being captured, exactly why this + /// counterpart exists (mirrors ComputeError() alongside + /// ComputeErrorAndEnergy() above). + virtual void ComputePrecisionWeightedError(float* e, const float* z, const float* mu, + const float* p, size_t batchSize, + size_t width) noexcept = 0; /// @brief y[b*width+i] += alpha * x[b*width+i] * factor[i], for every b /// in [0,batchSize), i in [0,width), AxpyInto with an added diff --git a/include/deepity/layers/ConvPCLayer.h b/include/deepity/layers/ConvPCLayer.h index 29df105..ca6bb13 100644 --- a/include/deepity/layers/ConvPCLayer.h +++ b/include/deepity/layers/ConvPCLayer.h @@ -70,7 +70,14 @@ class ConvPCLayer : public Layer { /// @brief Calculate energy/prediction errors for this layer. /// @return This layer's energy contribution at the current state. - float CalculateState() noexcept override; + float CalculateState() noexcept override { return CalculateState(true); } + /// @brief Same as CalculateState(), but can skip computing (and + /// returning) the energy. On CUDABackend, needEnergy=false also means + /// no blocking host sync, required for any settling-loop caller that + /// runs this inside a captured CUDA graph region (see + /// ConvPCNetwork::TrainStep()'s GPU branch). + /// @param needEnergy Whether to compute and return the energy. + float CalculateState(bool needEnergy) noexcept; /// @brief Update latent beliefs (z/r) via inference gradient. void UpdateState() noexcept override; /// @brief Hebbian/gradient weight update. diff --git a/include/deepity/layers/DiscriminativePCLayer.h b/include/deepity/layers/DiscriminativePCLayer.h index feb0539..988d513 100644 --- a/include/deepity/layers/DiscriminativePCLayer.h +++ b/include/deepity/layers/DiscriminativePCLayer.h @@ -83,7 +83,14 @@ namespace Deep /// E = \sum_l 1/2 p^{(l)} ||z^{(l)} - \mu^{(l)}||^2 - 1/2 \log p^{(l)} /// \f] /// @return This layer's energy contribution at the current state. - float CalculateState() noexcept override; + float CalculateState() noexcept override { return CalculateState(true); } + /// @brief Same as CalculateState(), but can skip computing (and + /// returning) the energy. On CUDABackend, needEnergy=false also + /// means no blocking host sync, required for any settling-loop + /// caller that runs this inside a captured CUDA graph region + /// (see DiscriminativePCNetwork::TrainStep()'s GPU branch). + /// @param needEnergy Whether to compute and return the energy. + float CalculateState(bool needEnergy) noexcept; /// @brief Computes the state derivatives for inference. /// diff --git a/include/deepity/layers/GaussSeidelPCLayer.h b/include/deepity/layers/GaussSeidelPCLayer.h index d056b1c..07a97fa 100644 --- a/include/deepity/layers/GaussSeidelPCLayer.h +++ b/include/deepity/layers/GaussSeidelPCLayer.h @@ -112,8 +112,14 @@ namespace Deep /// using this layer's own (just-updated) z as the target and /// layerBelow's FRESH mu (from its ComputePrediction() call) as /// the prediction. - /// @return This layer's energy contribution at the current state. - float ComputeError() noexcept; + /// @return This layer's energy contribution at the current state, + /// or 0.0f if needEnergy is false. + /// @param needEnergy Whether to compute and return the energy. On + /// CUDABackend, false also means no blocking host sync, required + /// for any settling-loop caller that runs this inside a captured + /// CUDA graph region (see GaussSeidelPCNetwork::TrainStep()'s GPU + /// branch). `e` itself is always correctly updated either way. + float ComputeError(bool needEnergy = true) noexcept; /// @brief Computes weight updates via gradient descent, with L2 /// weight decay. Called once after the full settling loop diff --git a/include/deepity/layers/SimpleConvPCLayer.h b/include/deepity/layers/SimpleConvPCLayer.h index a33607d..81251d5 100644 --- a/include/deepity/layers/SimpleConvPCLayer.h +++ b/include/deepity/layers/SimpleConvPCLayer.h @@ -57,7 +57,13 @@ namespace Deep ActivationType dType = ActivationType::dRELU, IComputeBackend *backend = nullptr); - float CalculateState() noexcept override; + float CalculateState() noexcept override { return CalculateState(true); } + /// @brief Same as CalculateState(), but can skip computing (and + /// returning) the energy. On CUDABackend, needEnergy=false also + /// means no blocking host sync, required for any settling-loop + /// caller that runs this inside a captured CUDA graph region. + /// @param needEnergy Whether to compute and return the energy. + float CalculateState(bool needEnergy) noexcept; void UpdateState() noexcept override; void UpdateWeights() noexcept override; diff --git a/include/deepity/networks/ConvPCNetwork.h b/include/deepity/networks/ConvPCNetwork.h index 0b83169..291cac8 100644 --- a/include/deepity/networks/ConvPCNetwork.h +++ b/include/deepity/networks/ConvPCNetwork.h @@ -62,7 +62,12 @@ namespace Deep void Clamp(const std::vector &input) noexcept; /// @brief Computes and returns the network's total energy at the /// current state, without changing it. - float CalculateState() noexcept; + /// @param needEnergy Whether to compute and return the energy. On + /// CUDABackend, false also means no blocking host sync anywhere + /// in the call, required when running this inside a captured + /// CUDA graph region (see TrainStep()'s GPU branch). Every + /// layer's own error (e) is always correctly updated either way. + float CalculateState(bool needEnergy = true) noexcept; /// @brief Runs one settling step on every layer. void UpdateState() noexcept; /// @brief Applies weight updates to every layer. diff --git a/include/deepity/networks/DiscriminativePCNetwork.h b/include/deepity/networks/DiscriminativePCNetwork.h index 67ca8a3..b176784 100644 --- a/include/deepity/networks/DiscriminativePCNetwork.h +++ b/include/deepity/networks/DiscriminativePCNetwork.h @@ -113,8 +113,12 @@ namespace Deep void Clamp(const std::vector &input); /// @brief Calculates the state of each layer. - /// @return Returns total energy - float CalculateState(); + /// @param needEnergy Whether to compute and return the energy. On + /// CUDABackend, false also means no blocking host sync anywhere + /// in the call, required when running this inside a captured + /// CUDA graph region. + /// @return Returns total energy, or 0.0f if needEnergy is false. + float CalculateState(bool needEnergy = true); /// @brief Updates each layer's state. void UpdateState(); diff --git a/include/deepity/networks/GaussSeidelPCNetwork.h b/include/deepity/networks/GaussSeidelPCNetwork.h index 0b87298..132d9f2 100644 --- a/include/deepity/networks/GaussSeidelPCNetwork.h +++ b/include/deepity/networks/GaussSeidelPCNetwork.h @@ -96,10 +96,15 @@ namespace Deep /// @brief Runs ONE full Gauss-Seidel settling step (all three /// sweeps, in order) across every layer. + /// @param needEnergy Whether to compute and return the energy. On + /// CUDABackend, false also means no blocking host sync anywhere + /// in the call, required when running this inside a captured + /// CUDA graph region. Sweep 1/2 (UpdateState/ComputePrediction) + /// and sweep 3's error update always run regardless. /// @return Total energy, summed from every layer's - /// ComputeError(). Meaningful only after all three sweeps have - /// run for this step. - float Step() noexcept; + /// ComputeError(), or 0.0f if needEnergy is false. Meaningful + /// only after all three sweeps have run for this step. + float Step(bool needEnergy = true) noexcept; /// @brief Updates every non-terminal layer's weights. Called /// once after the full settling loop completes. diff --git a/include/deepity/networks/SimpleConvPCNetwork.h b/include/deepity/networks/SimpleConvPCNetwork.h index e9e9890..cb83d0f 100644 --- a/include/deepity/networks/SimpleConvPCNetwork.h +++ b/include/deepity/networks/SimpleConvPCNetwork.h @@ -70,7 +70,11 @@ namespace Deep /// @brief Computes and returns the network's total energy at the /// current state, without changing it. - float CalculateState() noexcept; + /// @param needEnergy Whether to compute and return the energy. On + /// CUDABackend, false also means no blocking host sync anywhere + /// in the call, required when running this inside a captured + /// CUDA graph region. + float CalculateState(bool needEnergy = true) noexcept; /// @brief Runs one settling step on every layer. void UpdateState() noexcept; diff --git a/src/backend/CPUBackend.cpp b/src/backend/CPUBackend.cpp index 928fa22..e4ab235 100644 --- a/src/backend/CPUBackend.cpp +++ b/src/backend/CPUBackend.cpp @@ -128,6 +128,15 @@ void CPUBackend::SumRows(float* dst, const float* src, size_t batchSize, size_t cblas_saxpy(width, 1.0f, src + b * width, 1, dst, 1); } +float CPUBackend::Sum(const float* buf, size_t n) noexcept +{ + float total = 0.0f; +#pragma omp parallel for schedule(static) reduction(+ : total) if (n > 4096 && !omp_in_parallel()) + for (size_t i = 0; i < n; ++i) + total += buf[i]; + return total; +} + void CPUBackend::Scale(float* buf, size_t n, float alpha) noexcept { cblas_sscal(n, alpha, buf, 1); @@ -471,6 +480,16 @@ float CPUBackend::ComputePrecisionWeightedErrorAndEnergy(float* e, const float* return energy; } +void CPUBackend::ComputePrecisionWeightedError(float* e, const float* z, const float* mu, + const float* /*p*/, size_t batchSize, + size_t width) noexcept +{ + const size_t n = batchSize * width; +#pragma omp parallel for schedule(static) if (n > 256 && !omp_in_parallel()) + for (size_t idx = 0; idx < n; ++idx) + e[idx] = z[idx] - mu[idx]; +} + void CPUBackend::AxpyBroadcastInto(float* y, const float* x, const float* factor, size_t batchSize, size_t width, float alpha) noexcept { diff --git a/src/backend/cuda/CUDABackendFusedOps.cu b/src/backend/cuda/CUDABackendFusedOps.cu index 4694388..48c42f9 100644 --- a/src/backend/cuda/CUDABackendFusedOps.cu +++ b/src/backend/cuda/CUDABackendFusedOps.cu @@ -208,11 +208,19 @@ void CUDABackend::FusedActivationDerivativeMultiply(float* dst, const float* a, // a single device accumulator via atomicAdd. Simple and obviously correct // rather than a tree reduction, "port first, optimize" (see // RepackForBatchedGemm's own doc comment for the same reasoning applied -// elsewhere in this file), and this isn't inside any graph-captured -// region: the ConvPCLayer/DiscriminativePCLayer family that needs this -// doesn't use CUDA graph capture at all (unlike SimplePCNetwork/ -// FullPCNetwork/DirectKPPCNetwork), so a plain synchronous allocation -// here is safe. +// elsewhere in this file). +// +// IMPORTANT, read before touching this function: ConvPCNetwork and +// DiscriminativePCNetwork's TrainStep()/TrainStepWithProjection() DO now +// use CUDA graph capture (this comment previously said otherwise -- that +// was true when written, no longer is). This kernel launch itself is +// fine to capture, but CUDABackend::ComputePrecisionWeightedErrorAndEnergy() +// below it does a synchronous cudaMemcpyAsync+cudaStreamSynchronize to +// return the energy as a host float, which is illegal while a stream is +// being captured and crashes. Every call site inside a settling loop +// that's ever captured MUST use ComputePrecisionWeightedError() (below, +// no energy, no sync) instead, exactly like ComputeError() exists +// alongside ComputeErrorAndEnergy() for the same reason. __global__ void PrecisionWeightedErrorEnergyKernel(float* e, const float* z, const float* mu, const float* p, size_t n, size_t width, float* energyAccum) @@ -229,6 +237,15 @@ __global__ void PrecisionWeightedErrorEnergyKernel(float* e, const float* z, con } } +// No-energy counterpart: just e = z - mu, no atomicAdd/reduction/sync at +// all, safe to call from inside a captured region. +__global__ void PrecisionWeightedErrorKernel(float* e, const float* z, const float* mu, size_t n) +{ + size_t idx = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) + e[idx] = z[idx] - mu[idx]; +} + float CUDABackend::ComputePrecisionWeightedErrorAndEnergy(float* e, const float* z, const float* mu, const float* p, size_t batchSize, size_t width) noexcept @@ -256,6 +273,55 @@ float CUDABackend::ComputePrecisionWeightedErrorAndEnergy(float* e, const float* return total; } +void CUDABackend::ComputePrecisionWeightedError(float* e, const float* z, const float* mu, + const float* /*p*/, size_t batchSize, + size_t width) noexcept +{ + size_t n = batchSize * width; + if (!e || !z || !mu || n == 0) + return; + + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + PrecisionWeightedErrorKernel<<>>(e, z, mu, n); + CHECK_CUDA_LAUNCH(); +} + +// One thread per element, same atomicAdd-into-one-accumulator pattern as +// PrecisionWeightedErrorEnergyKernel above. +__global__ void SumKernel(const float* buf, size_t n, float* sumAccum) +{ + size_t idx = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) + atomicAdd(sumAccum, buf[idx]); +} + +float CUDABackend::Sum(const float* buf, size_t n) noexcept +{ + if (!buf || n == 0) + return 0.0f; + + float* sumAccum = Allocate(1); + if (!sumAccum) + return 0.0f; + cudaMemsetAsync(sumAccum, 0, sizeof(float), stream); + + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + SumKernel<<>>(buf, n, sumAccum); + CHECK_CUDA_LAUNCH(); + + // Synchronous readback, same as ComputeErrorAndEnergy and friends -- + // see this method's own IComputeBackend doc: never call from inside a + // captured graph region. + float total = 0.0f; + cudaMemcpyAsync(&total, sumAccum, sizeof(float), cudaMemcpyDeviceToHost, stream); + cudaStreamSynchronize(stream); + Free(sumAccum); + + return total; +} + __global__ void UpdatePrecisionFromErrorKernel(float* p, float* log_p, const float* e, size_t batchSize, size_t width, float pr) { diff --git a/src/layers/ConvPCLayer.cpp b/src/layers/ConvPCLayer.cpp index 72e5f1b..eaf1ad8 100644 --- a/src/layers/ConvPCLayer.cpp +++ b/src/layers/ConvPCLayer.cpp @@ -204,27 +204,37 @@ namespace Deep backend->RandomizeNormal(W, Wsz, 0.0f, limit, seed); } - float ConvPCLayer::CalculateState() noexcept + float ConvPCLayer::CalculateState(bool needEnergy) noexcept { size_t ownSize = (size_t)inChannels * inHeight * inWidth; size_t ownStateSize = (size_t)batchSize * ownSize; // Zero-Energy bypass, matches the original's -0.5*sum(log_p) // formula for the input layer (no e term at all, since there's - // no layerBelow to compare against). + // no layerBelow to compare against). backend->Sum() is a + // synchronous GPU readback (same constraint as + // ComputePrecisionWeightedErrorAndEnergy below), so it's only + // computed when actually needed -- also fixes a real bug this + // branch had before any needEnergy existed: indexing log_p[i] + // directly from host code is illegal once log_p is a GPU device + // pointer (DeviceMemoryArena), regardless of graph capture. if (layerBelow == nullptr) { backend->Zero(e, ownStateSize); float totalEnergy = 0.0f; - for (size_t i = 0; i < ownSize; ++i) - totalEnergy -= 0.5f * log_p[i] * batchSize; + if (needEnergy) + totalEnergy = -0.5f * (float)batchSize * backend->Sum(log_p, ownSize); if (outChannels > 0) ComputeMuOnly(); return totalEnergy; } - float totalEnergy = backend->ComputePrecisionWeightedErrorAndEnergy( - e, z, layerBelow->mu, p, batchSize, ownSize); + float totalEnergy = 0.0f; + if (needEnergy) + totalEnergy = backend->ComputePrecisionWeightedErrorAndEnergy( + e, z, layerBelow->mu, p, batchSize, ownSize); + else + backend->ComputePrecisionWeightedError(e, z, layerBelow->mu, p, batchSize, ownSize); if (outChannels > 0) ComputeMuOnly(); diff --git a/src/layers/DiscriminativePCLayer.cpp b/src/layers/DiscriminativePCLayer.cpp index 8921b83..96e3033 100644 --- a/src/layers/DiscriminativePCLayer.cpp +++ b/src/layers/DiscriminativePCLayer.cpp @@ -51,24 +51,35 @@ namespace Deep backend->RandomizeNormal(W, Wsz, 0.0f, limit, seed); } - float DiscriminativePCLayer::CalculateState() noexcept + float DiscriminativePCLayer::CalculateState(bool needEnergy) noexcept { const size_t N = (size_t)batchSize * size; + // See ConvPCLayer::CalculateState()'s identical comment: this + // branch indexed log_p[i] from host code unconditionally before, + // which is illegal once log_p is a GPU device pointer + // (DeviceMemoryArena) -- fixed by routing through backend->Sum() + // (a synchronous GPU readback, same constraint as + // ComputePrecisionWeightedErrorAndEnergy below, so only computed + // when actually needed). if (layerBelow == nullptr) { backend->Zero(e, N); float totalEnergy = 0.0f; - for (size_t i = 0; i < size; ++i) - totalEnergy -= 0.5f * log_p[i] * batchSize; + if (needEnergy) + totalEnergy = -0.5f * (float)batchSize * backend->Sum(log_p, size); if (nextSize > 0) ComputeMuOnly(); return totalEnergy; } - float totalEnergy = backend->ComputePrecisionWeightedErrorAndEnergy( - e, z, layerBelow->mu, p, batchSize, size); + float totalEnergy = 0.0f; + if (needEnergy) + totalEnergy = backend->ComputePrecisionWeightedErrorAndEnergy( + e, z, layerBelow->mu, p, batchSize, size); + else + backend->ComputePrecisionWeightedError(e, z, layerBelow->mu, p, batchSize, size); if (nextSize > 0) ComputeMuOnly(); diff --git a/src/layers/GaussSeidelPCLayer.cpp b/src/layers/GaussSeidelPCLayer.cpp index 1cbb67b..d3fce27 100644 --- a/src/layers/GaussSeidelPCLayer.cpp +++ b/src/layers/GaussSeidelPCLayer.cpp @@ -104,7 +104,7 @@ namespace Deep backend->AddBiasBroadcast(mu, b, batchSize, nextSize); } - float GaussSeidelPCLayer::ComputeError() noexcept + float GaussSeidelPCLayer::ComputeError(bool needEnergy) noexcept { size_t ownStateSize = (size_t)batchSize * size; @@ -114,7 +114,11 @@ namespace Deep return 0.0f; } - return backend->ComputeErrorAndEnergy(e, z, layerBelow->GetMu(), ownStateSize); + if (needEnergy) + return backend->ComputeErrorAndEnergy(e, z, layerBelow->GetMu(), ownStateSize); + + backend->ComputeError(e, z, layerBelow->GetMu(), ownStateSize); + return 0.0f; } void GaussSeidelPCLayer::UpdateWeights() noexcept diff --git a/src/layers/SimpleConvPCLayer.cpp b/src/layers/SimpleConvPCLayer.cpp index 2038bbd..46b4fd7 100644 --- a/src/layers/SimpleConvPCLayer.cpp +++ b/src/layers/SimpleConvPCLayer.cpp @@ -195,7 +195,7 @@ namespace Deep backend->RandomizeNormal(W, Wsz, 0.0f, limit, seed); } - float SimpleConvPCLayer::CalculateState() noexcept + float SimpleConvPCLayer::CalculateState(bool needEnergy) noexcept { size_t ownSize = (size_t)inChannels * inHeight * inWidth; size_t ownStateSize = (size_t)batchSize * ownSize; @@ -208,7 +208,11 @@ namespace Deep return 0.0f; } - float totalEnergy = backend->ComputeErrorAndEnergy(e, z, layerBelow->mu, ownStateSize); + float totalEnergy = 0.0f; + if (needEnergy) + totalEnergy = backend->ComputeErrorAndEnergy(e, z, layerBelow->mu, ownStateSize); + else + backend->ComputeError(e, z, layerBelow->mu, ownStateSize); if (outChannels > 0) ComputeMuOnly(); diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index c9c5d07..7829010 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -81,12 +81,12 @@ namespace Deep layers.front()->ClampState(input); } - float ConvPCNetwork::CalculateState() noexcept + float ConvPCNetwork::CalculateState(bool needEnergy) noexcept { float e = 0.0f; for (auto &l : layers) - e += l->CalculateState(); - return e; + e += l->CalculateState(needEnergy); + return needEnergy ? e : 0.0f; } void ConvPCNetwork::UpdateState() noexcept @@ -124,7 +124,11 @@ namespace Deep backend->BeginGraphCapture(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + // needEnergy=false: ComputePrecisionWeightedErrorAndEnergy's + // synchronous GPU readback is illegal while this stream is + // being captured. The real energy read happens in the + // fresh, uncaptured call at the end of this function. + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -149,7 +153,7 @@ namespace Deep { for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -159,7 +163,7 @@ namespace Deep { for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -242,7 +246,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -268,7 +272,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -279,7 +283,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); diff --git a/src/networks/DiscriminativePCNetwork.cpp b/src/networks/DiscriminativePCNetwork.cpp index 92c913d..2f26db9 100644 --- a/src/networks/DiscriminativePCNetwork.cpp +++ b/src/networks/DiscriminativePCNetwork.cpp @@ -74,12 +74,12 @@ namespace Deep layers.front()->ClampState(input); } - float DiscriminativePCNetwork::CalculateState() + float DiscriminativePCNetwork::CalculateState(bool needEnergy) { float e = 0.0f; for (auto &l : layers) - e += l->CalculateState(); - return e; + e += l->CalculateState(needEnergy); + return needEnergy ? e : 0.0f; } void DiscriminativePCNetwork::UpdateState() @@ -140,7 +140,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdatePrecision(); @@ -167,7 +167,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdatePrecision(); @@ -179,7 +179,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdatePrecision(); @@ -227,7 +227,7 @@ namespace Deep backend->BeginGraphCapture(); for (int t = 0; t < inferenceSteps; t++) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -252,7 +252,7 @@ namespace Deep { for (int t = 0; t < inferenceSteps; t++) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -262,7 +262,7 @@ namespace Deep { for (int t = 0; t < inferenceSteps; t++) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); diff --git a/src/networks/GaussSeidelPCNetwork.cpp b/src/networks/GaussSeidelPCNetwork.cpp index 575c101..d0d928d 100644 --- a/src/networks/GaussSeidelPCNetwork.cpp +++ b/src/networks/GaussSeidelPCNetwork.cpp @@ -47,7 +47,7 @@ namespace Deep layers.front()->ClampState(input); } - float GaussSeidelPCNetwork::Step() noexcept + float GaussSeidelPCNetwork::Step(bool needEnergy) noexcept { // Sweep 1: EVERY layer's z updates, using mu/e_above held over // from the end of the previous step. @@ -62,12 +62,16 @@ namespace Deep // Sweep 3: EVERY layer's error recomputes, using this step's // fresh z and layerBelow's fresh mu from sweep 2. Order among // layers doesn't matter here either. Energy is only meaningful - // starting from this point. + // starting from this point. needEnergy=false still correctly + // updates every layer's `e` (required for the next step's + // dynamics), just skips the energy value itself -- on + // CUDABackend that also means no blocking host sync, required + // when this runs inside a captured CUDA graph region. float totalEnergy = 0.0f; for (auto &l : layers) - totalEnergy += l->ComputeError(); + totalEnergy += l->ComputeError(needEnergy); - return totalEnergy; + return needEnergy ? totalEnergy : 0.0f; } void GaussSeidelPCNetwork::UpdateWeights() noexcept @@ -107,7 +111,7 @@ namespace Deep { backend->BeginGraphCapture(); for (int t = 0; t < inferenceSteps; ++t) - Step(); + Step(false); UpdateWeights(); bool captureOk = backend->EndGraphCapture(); @@ -129,14 +133,14 @@ namespace Deep else { for (int t = 0; t < inferenceSteps; ++t) - Step(); + Step(false); UpdateWeights(); } } else { for (int t = 0; t < inferenceSteps; ++t) - Step(); + Step(false); UpdateWeights(); } @@ -181,10 +185,13 @@ namespace Deep ProjectForward(); // Only the terminal layer's error is computed immediately // (e3 = z3 - mu3, target vs. projected prediction); hidden - // layers stay at zero. - GetTerminalLayer()->ComputeError(); + // layers stay at zero. needEnergy=false: this return value + // is discarded either way, and the call must not + // synchronously read back from the GPU while this stream + // is being captured. + GetTerminalLayer()->ComputeError(false); for (int t = 0; t < inferenceSteps; ++t) - Step(); + Step(false); UpdateWeights(); bool captureOk = backend->EndGraphCapture(); @@ -206,18 +213,18 @@ namespace Deep else { ProjectForward(); - GetTerminalLayer()->ComputeError(); + GetTerminalLayer()->ComputeError(false); for (int t = 0; t < inferenceSteps; ++t) - Step(); + Step(false); UpdateWeights(); } } else { ProjectForward(); - GetTerminalLayer()->ComputeError(); + GetTerminalLayer()->ComputeError(false); for (int t = 0; t < inferenceSteps; ++t) - Step(); + Step(false); UpdateWeights(); } diff --git a/src/networks/SimpleConvPCNetwork.cpp b/src/networks/SimpleConvPCNetwork.cpp index 841086e..bd9aef8 100644 --- a/src/networks/SimpleConvPCNetwork.cpp +++ b/src/networks/SimpleConvPCNetwork.cpp @@ -82,12 +82,12 @@ namespace Deep layers.front()->ClampState(input); } - float SimpleConvPCNetwork::CalculateState() noexcept + float SimpleConvPCNetwork::CalculateState(bool needEnergy) noexcept { float e = 0.0f; for (auto &l : layers) - e += l->CalculateState(); - return e; + e += l->CalculateState(needEnergy); + return needEnergy ? e : 0.0f; } void SimpleConvPCNetwork::UpdateState() noexcept @@ -115,7 +115,11 @@ namespace Deep backend->BeginGraphCapture(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + // needEnergy=false: the precision-free ComputeErrorAndEnergy + // still does a synchronous GPU readback, illegal while this + // stream is being captured. Real energy read happens in the + // fresh, uncaptured call at the end of this function. + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -140,7 +144,7 @@ namespace Deep { for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -150,7 +154,7 @@ namespace Deep { for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -221,7 +225,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -247,7 +251,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); @@ -258,7 +262,7 @@ namespace Deep ProjectForward(); for (int t = 0; t < inferenceSteps; ++t) { - CalculateState(); + CalculateState(false); UpdateState(); } UpdateWeights(); From 54de866bd2b7791e461a99c130564c4b4698ee21 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 01:14:10 -0400 Subject: [PATCH 03/12] Add explicit sync after every ReplayGraph() call (diagnostic/defensive 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. --- include/deepity/backend/CPUBackend.h | 2 ++ include/deepity/backend/CUDABackend.h | 2 ++ include/deepity/backend/IComputeBackend.h | 13 ++++++++++++- src/backend/cuda/CUDABackendCore.cu | 7 +++++++ src/networks/ConvPCNetwork.cpp | 8 ++++++++ src/networks/DirectKPPCNetwork.cpp | 1 + src/networks/DiscriminativePCNetwork.cpp | 2 ++ src/networks/FullConvPCNetwork.cpp | 1 + src/networks/FullPCNetwork.cpp | 6 ++++++ src/networks/GaussSeidelPCNetwork.cpp | 2 ++ src/networks/SimpleConvPCNetwork.cpp | 2 ++ src/networks/SimplePCNetwork.cpp | 1 + 12 files changed, 46 insertions(+), 1 deletion(-) diff --git a/include/deepity/backend/CPUBackend.h b/include/deepity/backend/CPUBackend.h index 3ce1364..911b273 100644 --- a/include/deepity/backend/CPUBackend.h +++ b/include/deepity/backend/CPUBackend.h @@ -23,6 +23,8 @@ class CPUBackend : public IComputeBackend } // nothing to fail on CPU /// @copydoc Deep::IComputeBackend::ReplayGraph void ReplayGraph() noexcept override {} + /// @copydoc Deep::IComputeBackend::Synchronize + void Synchronize() noexcept override {} /// @copydoc Deep::IComputeBackend::Allocate float* Allocate(size_t numFloats) override; diff --git a/include/deepity/backend/CUDABackend.h b/include/deepity/backend/CUDABackend.h index 0c761b2..baf6cc5 100644 --- a/include/deepity/backend/CUDABackend.h +++ b/include/deepity/backend/CUDABackend.h @@ -23,6 +23,8 @@ class CUDABackend : public IComputeBackend bool EndGraphCapture() noexcept override; /// @copydoc Deep::IComputeBackend::ReplayGraph void ReplayGraph() noexcept override; + /// @copydoc Deep::IComputeBackend::Synchronize + void Synchronize() noexcept override; /// @copydoc Deep::IComputeBackend::Allocate float* Allocate(size_t numFloats) override; diff --git a/include/deepity/backend/IComputeBackend.h b/include/deepity/backend/IComputeBackend.h index ba3ca43..38d52e8 100644 --- a/include/deepity/backend/IComputeBackend.h +++ b/include/deepity/backend/IComputeBackend.h @@ -31,8 +31,19 @@ class IComputeBackend virtual bool EndGraphCapture() noexcept = 0; /// @brief Launches a graph previously captured and instantiated by /// BeginGraphCapture()/EndGraphCapture(). No-op if EndGraphCapture() - /// never returned true. + /// never returned true. cudaGraphLaunch() is asynchronous: this + /// returns as soon as the launch is ENQUEUED, not once it completes. + /// Ordinarily that's fine (CUDA guarantees same-stream ordering, so a + /// later call that needs the result just needs to be queued on the + /// same stream too) -- but call Synchronize() afterward if you need + /// the replayed work to have definitely finished, e.g. before reading + /// results back on the host through any path that ISN'T itself a + /// synchronous backend call. virtual void ReplayGraph() noexcept = 0; + /// @brief Blocks until all previously enqueued work on this backend's + /// stream has completed. A no-op on CPUBackend (every CPU call is + /// already synchronous). + virtual void Synchronize() noexcept = 0; // Memory diff --git a/src/backend/cuda/CUDABackendCore.cu b/src/backend/cuda/CUDABackendCore.cu index b5dc368..ea974dc 100644 --- a/src/backend/cuda/CUDABackendCore.cu +++ b/src/backend/cuda/CUDABackendCore.cu @@ -90,6 +90,13 @@ void CUDABackend::ReplayGraph() noexcept std::cerr << "cudaGraphLaunch failed: " << cudaGetErrorString(err) << "\n"; } +void CUDABackend::Synchronize() noexcept +{ + cudaError_t err = cudaStreamSynchronize(stream); + if (err != cudaSuccess) + std::cerr << "cudaStreamSynchronize failed: " << cudaGetErrorString(err) << "\n"; +} + float* CUDABackend::Allocate(size_t numFloats) { float* ptr = nullptr; diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index 7829010..6365a1e 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -148,6 +148,13 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + // cudaGraphLaunch() only enqueues the replay; the fresh + // energy readout below needs it to have actually finished. + // Same-stream ordering should guarantee this without an + // explicit wait, but this costs nothing extra (the + // ...AndEnergy() call right after does its own blocking + // sync regardless) and removes any doubt. + backend->Synchronize(); } else { @@ -266,6 +273,7 @@ namespace Deep if (graphCapturedWithProjection) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/DirectKPPCNetwork.cpp b/src/networks/DirectKPPCNetwork.cpp index f48a26e..c0ffe5c 100644 --- a/src/networks/DirectKPPCNetwork.cpp +++ b/src/networks/DirectKPPCNetwork.cpp @@ -132,6 +132,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/DiscriminativePCNetwork.cpp b/src/networks/DiscriminativePCNetwork.cpp index 2f26db9..39e8247 100644 --- a/src/networks/DiscriminativePCNetwork.cpp +++ b/src/networks/DiscriminativePCNetwork.cpp @@ -161,6 +161,7 @@ namespace Deep if (graphCapturedWithProjection) { backend->ReplayGraph(); + backend->Synchronize(); } else { @@ -247,6 +248,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/FullConvPCNetwork.cpp b/src/networks/FullConvPCNetwork.cpp index 7b3df35..2019d1d 100644 --- a/src/networks/FullConvPCNetwork.cpp +++ b/src/networks/FullConvPCNetwork.cpp @@ -175,6 +175,7 @@ float FullConvPCNetwork::TrainStep(const std::vector& x, const std::vecto if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/FullPCNetwork.cpp b/src/networks/FullPCNetwork.cpp index ab46e37..c089355 100644 --- a/src/networks/FullPCNetwork.cpp +++ b/src/networks/FullPCNetwork.cpp @@ -170,7 +170,13 @@ float FullPCNetwork::TrainStep(const std::vector& x, const std::vectorReplayGraph(); + // cudaGraphLaunch() only enqueues the replay; see + // ConvPCNetwork::TrainStep()'s identical comment for why this + // matters even though same-stream ordering should already cover it. + backend->Synchronize(); + } else { ProjectForward(); diff --git a/src/networks/GaussSeidelPCNetwork.cpp b/src/networks/GaussSeidelPCNetwork.cpp index d0d928d..e046595 100644 --- a/src/networks/GaussSeidelPCNetwork.cpp +++ b/src/networks/GaussSeidelPCNetwork.cpp @@ -129,6 +129,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { @@ -209,6 +210,7 @@ namespace Deep if (graphCapturedWithProjection) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/SimpleConvPCNetwork.cpp b/src/networks/SimpleConvPCNetwork.cpp index bd9aef8..1a483d0 100644 --- a/src/networks/SimpleConvPCNetwork.cpp +++ b/src/networks/SimpleConvPCNetwork.cpp @@ -139,6 +139,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { @@ -245,6 +246,7 @@ namespace Deep if (graphCapturedWithProjection) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/SimplePCNetwork.cpp b/src/networks/SimplePCNetwork.cpp index 3889d2f..5ad2a96 100644 --- a/src/networks/SimplePCNetwork.cpp +++ b/src/networks/SimplePCNetwork.cpp @@ -174,6 +174,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { From 9ade9212a857e7d54d48cc3a88cbcdae9d8e59e8 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 01:24:01 -0400 Subject: [PATCH 04/12] Add temporary std::cerr instrumentation to ConvPCNetwork::TrainStep()'s 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. --- src/networks/ConvPCNetwork.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index 6365a1e..5ddc808 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -119,8 +119,13 @@ namespace Deep if (device == DeviceType::DEVICE_GPU) { + std::cerr << "[DIAG] TrainStep called: inferenceSteps=" << inferenceSteps + << " graphCaptured(before)=" << graphCaptured + << " capturedInferenceSteps(before)=" << capturedInferenceSteps << "\n"; + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) { + std::cerr << "[DIAG] entering (re)capture branch\n"; backend->BeginGraphCapture(); for (int t = 0; t < inferenceSteps; ++t) { @@ -133,6 +138,7 @@ namespace Deep } UpdateWeights(); bool captureOk = backend->EndGraphCapture(); + std::cerr << "[DIAG] EndGraphCapture() returned captureOk=" << captureOk << "\n"; if (captureOk) { @@ -143,10 +149,17 @@ namespace Deep { std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; } + std::cerr << "[DIAG] after capture: graphCaptured=" << graphCaptured + << " capturedInferenceSteps=" << capturedInferenceSteps << "\n"; + } + else + { + std::cerr << "[DIAG] skipping capture, reusing existing graph\n"; } if (graphCaptured) { + std::cerr << "[DIAG] calling ReplayGraph()\n"; backend->ReplayGraph(); // cudaGraphLaunch() only enqueues the replay; the fresh // energy readout below needs it to have actually finished. From b25958b14c5fc46cdf31385590c83e31d6b4b4ac Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 14:57:32 -0400 Subject: [PATCH 05/12] Add W[0] weight-buffer diagnostic around the recapture's ReplayGraph() 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. --- src/networks/ConvPCNetwork.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index 5ddc808..35745f6 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -119,9 +119,12 @@ namespace Deep if (device == DeviceType::DEVICE_GPU) { + float w0Before = 0.0f; + backend->CopyToHost(&w0Before, layers.front()->GetWeights(), 1); std::cerr << "[DIAG] TrainStep called: inferenceSteps=" << inferenceSteps << " graphCaptured(before)=" << graphCaptured - << " capturedInferenceSteps(before)=" << capturedInferenceSteps << "\n"; + << " capturedInferenceSteps(before)=" << capturedInferenceSteps + << " W[0](before)=" << w0Before << "\n"; if (!graphCaptured || capturedInferenceSteps != inferenceSteps) { @@ -168,6 +171,10 @@ namespace Deep // ...AndEnergy() call right after does its own blocking // sync regardless) and removes any doubt. backend->Synchronize(); + + float w0After = 0.0f; + backend->CopyToHost(&w0After, layers.front()->GetWeights(), 1); + std::cerr << "[DIAG] W[0](after replay+sync)=" << w0After << "\n"; } else { From d9ad3b3e1cd202594841f76e4cf32b7d63924c3f Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 15:06:16 -0400 Subject: [PATCH 06/12] Invalidate clamped layers' mu-cache after a (re)capture, before the fresh 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. --- include/deepity/layers/ConvPCLayer.h | 16 ++++++++++++++++ src/networks/ConvPCNetwork.cpp | 11 +++++++++++ 2 files changed, 27 insertions(+) diff --git a/include/deepity/layers/ConvPCLayer.h b/include/deepity/layers/ConvPCLayer.h index ca6bb13..2e47c10 100644 --- a/include/deepity/layers/ConvPCLayer.h +++ b/include/deepity/layers/ConvPCLayer.h @@ -172,6 +172,22 @@ class ConvPCLayer : public Layer { /// @brief Whether ClampState() is currently active on this layer. bool IsClamped() const noexcept { return isClamped; } + /// @brief Forces the next ComputeMuOnly() call to recompute instead of + /// reusing cachedMu, without otherwise touching clamp state. + /// + /// ClampState() normally does this (muCacheValid = false), and runs + /// fresh at the top of every TrainStep() call. But a CUDA-graph (re)capture + /// replays ClampState() only implicitly, by re-executing the whole + /// settling loop's C++ as part of recording the graph: that loop's last + /// ComputeMuOnly() call leaves muCacheValid = true as a side effect, and + /// nothing clears it again before the graph is replayed. The following + /// uncaptured, fresh CalculateState() read (used for the returned/printed + /// energy) then wrongly takes the cached branch, returning mu from + /// BEFORE this call's weight update instead of after it. Callers that + /// just finished a capture should invalidate the cache here before that + /// read, exactly like a fresh ClampState() call would have. + void InvalidateMuCache() noexcept { muCacheValid = false; } + /// @brief Sets the layer immediately above this one in the network. /// @param above Pointer to the layer above; may be nullptr for a /// terminal layer. diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index 35745f6..d560ed4 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -147,6 +147,17 @@ namespace Deep { graphCaptured = true; capturedInferenceSteps = inferenceSteps; + + // The settling loop just recorded above ran as real C++, + // leaving muCacheValid=true whenever it last touched a + // clamped layer. Nothing will reset that before the fresh + // CalculateState() read below, so it would wrongly reuse + // cachedMu (computed from this call's PRE-update weights, + // the last time the loop's recording touched it) instead + // of recomputing from the weights ReplayGraph() is about + // to produce. See ConvPCLayer::InvalidateMuCache(). + for (auto &l : layers) + l->InvalidateMuCache(); } else { From 4297bf174fd4dfa4b2b4c403731ec47367bbc22c Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 15:08:59 -0400 Subject: [PATCH 07/12] Remove temporary [DIAG] instrumentation now that the stale-energy bug 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. --- src/networks/ConvPCNetwork.cpp | 20 -------------------- 1 file changed, 20 deletions(-) diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index d560ed4..e9ef042 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -119,16 +119,8 @@ namespace Deep if (device == DeviceType::DEVICE_GPU) { - float w0Before = 0.0f; - backend->CopyToHost(&w0Before, layers.front()->GetWeights(), 1); - std::cerr << "[DIAG] TrainStep called: inferenceSteps=" << inferenceSteps - << " graphCaptured(before)=" << graphCaptured - << " capturedInferenceSteps(before)=" << capturedInferenceSteps - << " W[0](before)=" << w0Before << "\n"; - if (!graphCaptured || capturedInferenceSteps != inferenceSteps) { - std::cerr << "[DIAG] entering (re)capture branch\n"; backend->BeginGraphCapture(); for (int t = 0; t < inferenceSteps; ++t) { @@ -141,7 +133,6 @@ namespace Deep } UpdateWeights(); bool captureOk = backend->EndGraphCapture(); - std::cerr << "[DIAG] EndGraphCapture() returned captureOk=" << captureOk << "\n"; if (captureOk) { @@ -163,17 +154,10 @@ namespace Deep { std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; } - std::cerr << "[DIAG] after capture: graphCaptured=" << graphCaptured - << " capturedInferenceSteps=" << capturedInferenceSteps << "\n"; - } - else - { - std::cerr << "[DIAG] skipping capture, reusing existing graph\n"; } if (graphCaptured) { - std::cerr << "[DIAG] calling ReplayGraph()\n"; backend->ReplayGraph(); // cudaGraphLaunch() only enqueues the replay; the fresh // energy readout below needs it to have actually finished. @@ -182,10 +166,6 @@ namespace Deep // ...AndEnergy() call right after does its own blocking // sync regardless) and removes any doubt. backend->Synchronize(); - - float w0After = 0.0f; - backend->CopyToHost(&w0After, layers.front()->GetWeights(), 1); - std::cerr << "[DIAG] W[0](after replay+sync)=" << w0After << "\n"; } else { From 287ef8965524cd150b269e043a0fe81d0b11e4cf Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 15:13:26 -0400 Subject: [PATCH 08/12] Fix the same stale-mu-cache-after-recapture bug in every other graph-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. --- include/deepity/layers/DirectKPPCLayer.h | 8 ++++++++ .../deepity/layers/DiscriminativePCLayer.h | 9 +++++++++ include/deepity/layers/SimpleConvPCLayer.h | 9 +++++++++ include/deepity/layers/SimplePCLayer.h | 9 +++++++++ src/networks/ConvPCNetwork.cpp | 7 +++++++ src/networks/DirectKPPCNetwork.cpp | 10 ++++++++++ src/networks/DiscriminativePCNetwork.cpp | 20 +++++++++++++++++++ src/networks/FullConvPCNetwork.cpp | 8 ++++++++ src/networks/FullPCNetwork.cpp | 8 ++++++++ src/networks/SimpleConvPCNetwork.cpp | 17 ++++++++++++++++ src/networks/SimplePCNetwork.cpp | 10 ++++++++++ 11 files changed, 115 insertions(+) diff --git a/include/deepity/layers/DirectKPPCLayer.h b/include/deepity/layers/DirectKPPCLayer.h index 2dc3723..191c688 100644 --- a/include/deepity/layers/DirectKPPCLayer.h +++ b/include/deepity/layers/DirectKPPCLayer.h @@ -173,6 +173,14 @@ class DirectKPPCLayer : public Layer { /// first time by moving ProjectForward() inside graph capture. bool IsClamped() const noexcept { return isClamped; } + /// @brief Forces the next ComputeMuOnly() call to recompute mu instead + /// of reusing a cached value, without otherwise touching clamp state. + /// Needed after a CUDA-graph (re)capture: the recording pass runs the + /// settling loop as real C++, leaving muCacheValid=true as a side + /// effect, which would otherwise make the following fresh, uncaptured + /// energy read reuse mu from before this call's own weight update. + void InvalidateMuCache() noexcept { muCacheValid = false; } + /// @brief Sets the optimizer for W. Call before Compile(). void SetOptimizer(OptimizerType o) noexcept { opt = o; } /// @brief Sets the optimizer for Psi. Call before Compile(). diff --git a/include/deepity/layers/DiscriminativePCLayer.h b/include/deepity/layers/DiscriminativePCLayer.h index 988d513..ec8d99f 100644 --- a/include/deepity/layers/DiscriminativePCLayer.h +++ b/include/deepity/layers/DiscriminativePCLayer.h @@ -126,6 +126,15 @@ namespace Deep /// @brief Whether ClampState() is currently active on this layer. bool IsClamped() const noexcept { return isClamped; } + /// @brief Forces the next ComputeMuOnly() call to recompute mu + /// instead of reusing a cached value, without otherwise touching + /// clamp state. Needed after a CUDA-graph (re)capture: the + /// recording pass runs the settling loop as real C++, leaving + /// muCacheValid=true as a side effect, which would otherwise make + /// the following fresh, uncaptured energy read reuse mu from + /// before this call's own weight update. + void InvalidateMuCache() noexcept { muCacheValid = false; } + /// @brief Returns beliefs. /// @return float *z float *GetBeliefs() noexcept override { return z; } diff --git a/include/deepity/layers/SimpleConvPCLayer.h b/include/deepity/layers/SimpleConvPCLayer.h index 81251d5..1308ef4 100644 --- a/include/deepity/layers/SimpleConvPCLayer.h +++ b/include/deepity/layers/SimpleConvPCLayer.h @@ -111,6 +111,15 @@ namespace Deep /// @brief Whether ClampState() is currently active on this layer. bool IsClamped() const noexcept { return isClamped; } + /// @brief Forces the next ComputeMuOnly() call to recompute mu + /// instead of reusing a cached value, without otherwise touching + /// clamp state. Needed after a CUDA-graph (re)capture: the + /// recording pass runs the settling loop as real C++, leaving + /// muCacheValid=true as a side effect, which would otherwise make + /// the following fresh, uncaptured energy read reuse mu from + /// before this call's own weight update. + void InvalidateMuCache() noexcept { muCacheValid = false; } + /// @brief Sets the layer immediately above this one in the network. void SetLayerAbove(SimpleConvPCLayer *above) noexcept { layerAbove = above; } /// @brief Sets the layer immediately below this one in the network. diff --git a/include/deepity/layers/SimplePCLayer.h b/include/deepity/layers/SimplePCLayer.h index 0858f16..ebcff04 100644 --- a/include/deepity/layers/SimplePCLayer.h +++ b/include/deepity/layers/SimplePCLayer.h @@ -127,6 +127,15 @@ namespace Deep /// @brief Whether ClampState() is currently active on this layer. bool IsClamped() const noexcept { return isClamped; } + /// @brief Forces the next ComputeMuOnly() call to recompute mu + /// instead of reusing a cached value, without otherwise touching + /// clamp state. Needed after a CUDA-graph (re)capture: the + /// recording pass runs the settling loop as real C++, leaving + /// muCacheValid=true as a side effect, which would otherwise make + /// the following fresh, uncaptured energy read reuse mu from + /// before this call's own weight update. + void InvalidateMuCache() noexcept { muCacheValid = false; } + /// @brief Returns this layer's belief buffer. /// @return Pointer to this layer's `size`-length beliefs. float *GetBeliefs() noexcept override { return z; } diff --git a/src/networks/ConvPCNetwork.cpp b/src/networks/ConvPCNetwork.cpp index e9ef042..2e71230 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -274,6 +274,13 @@ namespace Deep { graphCapturedWithProjection = true; capturedInferenceStepsWithProjection = inferenceSteps; + + // See the identical fix in TrainStep() above: the + // recording pass above just ran this settling loop as + // real C++, leaving muCacheValid=true behind for any + // clamped layer it touched last. + for (auto &l : layers) + l->InvalidateMuCache(); } else { diff --git a/src/networks/DirectKPPCNetwork.cpp b/src/networks/DirectKPPCNetwork.cpp index c0ffe5c..c02cfb4 100644 --- a/src/networks/DirectKPPCNetwork.cpp +++ b/src/networks/DirectKPPCNetwork.cpp @@ -122,6 +122,16 @@ namespace Deep { graphCaptured = true; capturedInferenceSteps = inferenceSteps; + + // The recording pass above just ran this settling loop + // as real C++, leaving muCacheValid=true behind for any + // clamped layer it touched last; nothing resets that + // before the fresh energy read below, which would + // otherwise wrongly reuse mu from before this call's + // own weight update (see ConvPCNetwork::TrainStep()'s + // identical bug). + for (auto &l : layers) + l->InvalidateMuCache(); } else { diff --git a/src/networks/DiscriminativePCNetwork.cpp b/src/networks/DiscriminativePCNetwork.cpp index 39e8247..8c643d3 100644 --- a/src/networks/DiscriminativePCNetwork.cpp +++ b/src/networks/DiscriminativePCNetwork.cpp @@ -151,6 +151,16 @@ namespace Deep { graphCapturedWithProjection = true; capturedInferenceStepsWithProjection = inferenceSteps; + + // The recording pass above just ran this settling loop + // as real C++, leaving muCacheValid=true behind for any + // clamped layer it touched last; nothing resets that + // before the fresh energy read below, which would + // otherwise wrongly reuse mu from before this call's + // own weight update (see ConvPCNetwork::TrainStep()'s + // identical bug). + for (auto &l : layers) + l->InvalidateMuCache(); } else { @@ -238,6 +248,16 @@ namespace Deep { graphCaptured = true; capturedInferenceSteps = inferenceSteps; + + // The recording pass above just ran this settling loop + // as real C++, leaving muCacheValid=true behind for any + // clamped layer it touched last; nothing resets that + // before the fresh energy read below, which would + // otherwise wrongly reuse mu from before this call's + // own weight update (see ConvPCNetwork::TrainStep()'s + // identical bug). + for (auto &l : layers) + l->InvalidateMuCache(); } else { diff --git a/src/networks/FullConvPCNetwork.cpp b/src/networks/FullConvPCNetwork.cpp index 2019d1d..a382464 100644 --- a/src/networks/FullConvPCNetwork.cpp +++ b/src/networks/FullConvPCNetwork.cpp @@ -165,6 +165,14 @@ float FullConvPCNetwork::TrainStep(const std::vector& x, const std::vecto { graphCaptured = true; capturedInferenceSteps = inferenceSteps; + + // settleStep()'s own InvalidateMuCache() calls above only run + // under useIPC; without it, nothing clears a clamped layer's + // mu-cache after this recording pass, so the fresh CalculateState() + // read below would wrongly reuse mu from before this call's own + // weight update (see ConvPCNetwork::TrainStep()'s identical bug). + for (auto& l : layers) + l->InvalidateMuCache(); } else { diff --git a/src/networks/FullPCNetwork.cpp b/src/networks/FullPCNetwork.cpp index c089355..9771947 100644 --- a/src/networks/FullPCNetwork.cpp +++ b/src/networks/FullPCNetwork.cpp @@ -164,6 +164,14 @@ float FullPCNetwork::TrainStep(const std::vector& x, const std::vectorInvalidateMuCache(); } else std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; diff --git a/src/networks/SimpleConvPCNetwork.cpp b/src/networks/SimpleConvPCNetwork.cpp index 1a483d0..1f55e6e 100644 --- a/src/networks/SimpleConvPCNetwork.cpp +++ b/src/networks/SimpleConvPCNetwork.cpp @@ -129,6 +129,16 @@ namespace Deep { graphCaptured = true; capturedInferenceSteps = inferenceSteps; + + // The recording pass above just ran this settling loop + // as real C++, leaving muCacheValid=true behind for any + // clamped layer it touched last; nothing resets that + // before the fresh energy read below, which would + // otherwise wrongly reuse mu from before this call's + // own weight update (see ConvPCNetwork::TrainStep()'s + // identical bug). + for (auto &l : layers) + l->InvalidateMuCache(); } else { @@ -236,6 +246,13 @@ namespace Deep { graphCapturedWithProjection = true; capturedInferenceStepsWithProjection = inferenceSteps; + + // See the identical fix above in TrainStep(): the + // recording pass just ran this settling loop as real + // C++, leaving muCacheValid=true behind for any + // clamped layer it touched last. + for (auto &l : layers) + l->InvalidateMuCache(); } else { diff --git a/src/networks/SimplePCNetwork.cpp b/src/networks/SimplePCNetwork.cpp index 5ad2a96..a0c0c34 100644 --- a/src/networks/SimplePCNetwork.cpp +++ b/src/networks/SimplePCNetwork.cpp @@ -164,6 +164,16 @@ namespace Deep { graphCaptured = true; capturedInferenceSteps = inferenceSteps; + + // The recording pass above just ran this settling loop + // as real C++, leaving muCacheValid=true behind for any + // clamped layer it touched last; nothing resets that + // before the fresh energy read below, which would + // otherwise wrongly reuse mu from before this call's + // own weight update (see ConvPCNetwork::TrainStep()'s + // identical bug). + for (auto &l : layers) + l->InvalidateMuCache(); } else { From c6149874487ff85deabc8ffdeb2b9e37ec817328 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 15:46:52 -0400 Subject: [PATCH 09/12] Fixed pyright and clang errors in workflow. --- pydeepity/ConvolutionalPCN.py | 3 +-- pydeepity/__init__.py | 3 +++ pydeepity/_backend.py | 9 ++++++++- temp.py | 2 +- tests/tGelu.cpp | 5 +++-- tests/tReadme.cpp | 12 ++++++------ 6 files changed, 22 insertions(+), 12 deletions(-) diff --git a/pydeepity/ConvolutionalPCN.py b/pydeepity/ConvolutionalPCN.py index 423ef44..31d8544 100644 --- a/pydeepity/ConvolutionalPCN.py +++ b/pydeepity/ConvolutionalPCN.py @@ -243,8 +243,7 @@ def fit( if initial_lr is None: if self._learning_rate is None: raise RuntimeError("Learning rate has not been configured.") - - initial_lr = self._learning_rate + initial_lr = self._learning_rate _fit_with_progress( self, diff --git a/pydeepity/__init__.py b/pydeepity/__init__.py index 1615cf3..c5f5229 100644 --- a/pydeepity/__init__.py +++ b/pydeepity/__init__.py @@ -62,6 +62,9 @@ "GeLU", "Convolution", + # Backend + "dy", + # Backend utilities "StreamAlignedBatcher", "get_l2_cache_bytes", diff --git a/pydeepity/_backend.py b/pydeepity/_backend.py index 33f2bdc..2c9f8eb 100644 --- a/pydeepity/_backend.py +++ b/pydeepity/_backend.py @@ -1,5 +1,12 @@ try: - from . import pydeepity as dy + # Absolute form, not `from . import pydeepity as dy`: pyright can't + # resolve a relative import of a submodule that shares its name + # with the enclosing package (no stub exists for the compiled + # extension either way, but the relative form fails to resolve + # at all, which cascades into "dy is unknown import symbol" + # everywhere dy gets re-exported, outside this package's own + # reportAttributeAccessIssue-suppressed executionEnvironment). + import pydeepity.pydeepity as dy except ImportError as e: raise ImportError( "Could not load the compiled Deepity C++ backend.\nEnsure the package was installed correctly or compiled for your architecture." diff --git a/temp.py b/temp.py index c8931e0..cbe8101 100644 --- a/temp.py +++ b/temp.py @@ -19,7 +19,7 @@ def load_full_mnist(): } data_dir = "./data" os.makedirs(data_dir, exist_ok=True) - paths = {} + paths: dict[str, str] = {} for key, fname in files.items(): filepath = os.path.join(data_dir, fname) paths[key] = filepath diff --git a/tests/tGelu.cpp b/tests/tGelu.cpp index be2586d..b068069 100644 --- a/tests/tGelu.cpp +++ b/tests/tGelu.cpp @@ -15,6 +15,7 @@ * all produce errors far larger than this). */ #include +#include #include #include #include @@ -64,7 +65,7 @@ int main() // FMA structure) wouldn't show up in the n=1 scalar-only test above. printf("\n=== SIMD path (buffer large enough to exercise vectorized loop) ===\n"); const size_t simdSize = 64; // large enough for any SIMD width (16/8/4) with a real tail too - float *buf = static_cast(std::aligned_alloc(64, simdSize * sizeof(float))); + float *buf = static_cast(detail::AlignedAllocPortable(64, simdSize * sizeof(float))); for (size_t i = 0; i < simdSize; ++i) buf[i] = cases[i % cases.size()].x; @@ -80,7 +81,7 @@ int main() printf(" buf[%zu]: got %.6f, expected %.6f, err=%.6f FAIL\n", i, buf[i], expected, err); } printf(" (silent = all %zu SIMD-path values passed)\n", simdSize); - std::free(buf); + detail::AlignedFreePortable(buf); printf("\n%s\n", allPassed ? "PASS: gelu() matches true GELU within tolerance." : "FAIL: gelu() does not match expected values."); return allPassed ? 0 : 1; diff --git a/tests/tReadme.cpp b/tests/tReadme.cpp index 947b795..101d4c7 100644 --- a/tests/tReadme.cpp +++ b/tests/tReadme.cpp @@ -158,7 +158,7 @@ static void BM_Activation_DeepityTanh(benchmark::State &state) { size_t size = state.range(0); size_t bytes = (size * sizeof(float) + 63) & ~63; // 64-byte alignment padding - float *data = static_cast(std::aligned_alloc(64, bytes)); + float *data = static_cast(Deep::detail::AlignedAllocPortable(64, bytes)); std::fill_n(data, size, 0.5f); for (auto _ : state) @@ -166,7 +166,7 @@ static void BM_Activation_DeepityTanh(benchmark::State &state) Deep::tanh(data, size); benchmark::ClobberMemory(); } - std::free(data); + Deep::detail::AlignedFreePortable(data); } BENCHMARK(BM_Activation_DeepityTanh)->Arg(10048)->Arg(1000064); @@ -174,7 +174,7 @@ static void BM_Activation_DeepityDTanh(benchmark::State &state) { size_t size = state.range(0); size_t bytes = (size * sizeof(float) + 63) & ~63; - float *data = static_cast(std::aligned_alloc(64, bytes)); + float *data = static_cast(Deep::detail::AlignedAllocPortable(64, bytes)); std::fill_n(data, size, 0.5f); for (auto _ : state) @@ -182,7 +182,7 @@ static void BM_Activation_DeepityDTanh(benchmark::State &state) Deep::dTanh(data, size, true); benchmark::ClobberMemory(); } - std::free(data); + Deep::detail::AlignedFreePortable(data); } BENCHMARK(BM_Activation_DeepityDTanh)->Arg(10048)->Arg(1000064); @@ -202,7 +202,7 @@ static void BM_Activation_DeepitySigmoid(benchmark::State &state) { size_t size = state.range(0); size_t bytes = (size * sizeof(float) + 63) & ~63; - float *data = static_cast(std::aligned_alloc(64, bytes)); + float *data = static_cast(Deep::detail::AlignedAllocPortable(64, bytes)); std::fill_n(data, size, 0.5f); for (auto _ : state) @@ -210,7 +210,7 @@ static void BM_Activation_DeepitySigmoid(benchmark::State &state) Deep::sigmoid(data, size); benchmark::ClobberMemory(); } - std::free(data); + Deep::detail::AlignedFreePortable(data); } BENCHMARK(BM_Activation_DeepitySigmoid)->Arg(10048)->Arg(1000064); From 1e3786cd843d8ef137010c8052f6518bdfba90ca Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sat, 3 Oct 2026 16:03:08 -0400 Subject: [PATCH 10/12] Fixed pyright and clang errors in workflow x2. --- imagenet.py | 2 +- pydeepity/DKPPCN.py | 2 ++ pydeepity/FullPCN.py | 2 ++ pydeepity/GaussSeidelPCN.py | 2 ++ pydeepity/SequentialPCN.py | 2 ++ pydeepity/SimpleConvolutionalPCN.py | 2 ++ pydeepity/SimplePCN.py | 2 ++ pydeepity/_backend.py | 9 +-------- 8 files changed, 14 insertions(+), 9 deletions(-) diff --git a/imagenet.py b/imagenet.py index cb8ac99..72464a7 100644 --- a/imagenet.py +++ b/imagenet.py @@ -2,7 +2,7 @@ import os import sys from time import perf_counter -from pydeepity import dy +from pydeepity import dy # pyright: ignore[reportAttributeAccessIssue] -- dy is the compiled native extension; see pydeepity/_backend.py and pyrightconfig.json's pydeepity/bindings suppressions for the same reason from PIL import Image IMG_SIZE = 64 diff --git a/pydeepity/DKPPCN.py b/pydeepity/DKPPCN.py index 3fd8a2a..198bd43 100644 --- a/pydeepity/DKPPCN.py +++ b/pydeepity/DKPPCN.py @@ -1,3 +1,5 @@ +from __future__ import annotations + from typing import Optional import numpy as np diff --git a/pydeepity/FullPCN.py b/pydeepity/FullPCN.py index f6c7eae..ad7609b 100644 --- a/pydeepity/FullPCN.py +++ b/pydeepity/FullPCN.py @@ -1,3 +1,5 @@ +from __future__ import annotations + from typing import Optional import numpy as np diff --git a/pydeepity/GaussSeidelPCN.py b/pydeepity/GaussSeidelPCN.py index ee435ac..8f1ff31 100644 --- a/pydeepity/GaussSeidelPCN.py +++ b/pydeepity/GaussSeidelPCN.py @@ -1,3 +1,5 @@ +from __future__ import annotations + from typing import Optional import numpy as np diff --git a/pydeepity/SequentialPCN.py b/pydeepity/SequentialPCN.py index bcb8eb3..a8a33f8 100644 --- a/pydeepity/SequentialPCN.py +++ b/pydeepity/SequentialPCN.py @@ -1,3 +1,5 @@ +from __future__ import annotations + from ._backend import dy from .layer import Layer, Linear, Activation from .utils import _fit_with_progress diff --git a/pydeepity/SimpleConvolutionalPCN.py b/pydeepity/SimpleConvolutionalPCN.py index 81a7327..bc2139a 100644 --- a/pydeepity/SimpleConvolutionalPCN.py +++ b/pydeepity/SimpleConvolutionalPCN.py @@ -1,3 +1,5 @@ +from __future__ import annotations + from ._backend import dy from .layer import Layer, Convolution, Activation from .utils import _fit_with_progress diff --git a/pydeepity/SimplePCN.py b/pydeepity/SimplePCN.py index 4639b22..ba19d88 100644 --- a/pydeepity/SimplePCN.py +++ b/pydeepity/SimplePCN.py @@ -1,3 +1,5 @@ +from __future__ import annotations + from typing import Optional import numpy as np diff --git a/pydeepity/_backend.py b/pydeepity/_backend.py index 2c9f8eb..33f2bdc 100644 --- a/pydeepity/_backend.py +++ b/pydeepity/_backend.py @@ -1,12 +1,5 @@ try: - # Absolute form, not `from . import pydeepity as dy`: pyright can't - # resolve a relative import of a submodule that shares its name - # with the enclosing package (no stub exists for the compiled - # extension either way, but the relative form fails to resolve - # at all, which cascades into "dy is unknown import symbol" - # everywhere dy gets re-exported, outside this package's own - # reportAttributeAccessIssue-suppressed executionEnvironment). - import pydeepity.pydeepity as dy + from . import pydeepity as dy except ImportError as e: raise ImportError( "Could not load the compiled Deepity C++ backend.\nEnsure the package was installed correctly or compiled for your architecture." From c0eb0a45292ace18d5e82d03ca21a2e67cb38c80 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sun, 4 Oct 2026 12:39:12 -0400 Subject: [PATCH 11/12] Greater CUTLASS integration in cudaMalloc/free and activations --- include/deepity/backend/CUDABackend.h | 9 ++ include/deepity/backend/IComputeBackend.h | 29 ++-- src/backend/cuda/CUDABackendCore.cu | 3 + src/backend/cuda/CUDABackendFusedOps.cu | 104 +++++++++---- src/backend/cuda/CUDABackendGemm.cu | 177 ++++++++++++++++++++++ 5 files changed, 280 insertions(+), 42 deletions(-) diff --git a/include/deepity/backend/CUDABackend.h b/include/deepity/backend/CUDABackend.h index baf6cc5..7757227 100644 --- a/include/deepity/backend/CUDABackend.h +++ b/include/deepity/backend/CUDABackend.h @@ -174,5 +174,14 @@ class CUDABackend : public IComputeBackend float* onesVector = nullptr; size_t onesCapacity = 0; float* workspace = nullptr; + + /// @brief Single-float device scratch buffer for scalar reductions + /// (ComputePrecisionWeightedErrorAndEnergy(), Sum()). Allocated once in + /// the constructor instead of per-call: cudaMalloc/cudaFree are + /// synchronizing driver calls that stall the whole pipeline, and both + /// callers already block on cudaStreamSynchronize() before returning, + /// so reusing one persistent buffer across calls is safe -- no two + /// reductions on this backend instance can ever be in flight at once. + float* scalarScratch = nullptr; }; } // namespace Deep \ No newline at end of file diff --git a/include/deepity/backend/IComputeBackend.h b/include/deepity/backend/IComputeBackend.h index 38d52e8..677d58e 100644 --- a/include/deepity/backend/IComputeBackend.h +++ b/include/deepity/backend/IComputeBackend.h @@ -217,15 +217,26 @@ class IComputeBackend size_t batchSize, size_t nextSize) noexcept = 0; /// @brief Attempts a fused forward pass (GEMM + bias + activation in - /// one kernel, via CUTLASS on GPU) for RELU or LINEAR activation types - /// only, see CUDABackend's implementation for why other activation - /// types aren't attempted yet. Returns true if the fused path was used - /// (mu is fully computed, including bias and activation); false if the - /// caller should fall back to the existing MatMul+AddBiasBroadcast+ - /// Activation sequence (always false on CPUBackend, CPU has no fused - /// path, this is a GPU-only optimization). - /// @param actType Activation to fuse in; only RELU and LINEAR take the - /// fused path, everything else returns false immediately. + /// one kernel, via CUTLASS on GPU) for RELU, LINEAR, GELU, TANH, or + /// SIGMOID activation types; anything else returns false immediately. + /// Returns true if the fused path was used (mu is fully computed, + /// including bias and activation); false if the caller should fall + /// back to the existing MatMul+AddBiasBroadcast+Activation sequence + /// (always false on CPUBackend, CPU has no fused path, this is a + /// GPU-only optimization). + /// @warning As of this writing, every existing call site + /// (FullPCLayer::ComputeMuOnly(), DirectKPPCLayer::ComputeMuOnly()) + /// passes actType=LINEAR unconditionally -- their real activation is + /// applied via a separate ActivationInto() pass on the INPUT before + /// this GEMM runs, not fused into its epilogue. The RELU/GELU/TANH/ + /// SIGMOID branches below are real and GPU-tested-correct in + /// isolation, but exercising them for an actual layer requires moving + /// that activation into this call instead (see CUDABackendGemm.cu's + /// file comment before doing so: it's a data-flow change to a settling + /// loop, not just a new branch). + /// @param actType Activation to fuse in; RELU, LINEAR, GELU, TANH, and + /// SIGMOID take the fused path, everything else returns false + /// immediately. /// @param zF Activated input, shape [batchSize, size], row-major. /// @param W Weight matrix, shape [nextSize, size], row-major. /// @param bias Bias vector, shape [nextSize]. diff --git a/src/backend/cuda/CUDABackendCore.cu b/src/backend/cuda/CUDABackendCore.cu index ea974dc..cb39f1d 100644 --- a/src/backend/cuda/CUDABackendCore.cu +++ b/src/backend/cuda/CUDABackendCore.cu @@ -23,6 +23,7 @@ CUDABackend::CUDABackend() cudaStreamCreate(&this->stream); cublasCreate(&this->handle); cublasSetStream(this->handle, this->stream); + cudaMalloc(&scalarScratch, sizeof(float)); } CUDABackend::~CUDABackend() @@ -36,6 +37,8 @@ CUDABackend::~CUDABackend() cudaFree(onesVector); if (workspace) cudaFree(workspace); + if (scalarScratch) + cudaFree(scalarScratch); cublasDestroy(this->handle); cudaStreamDestroy(this->stream); } diff --git a/src/backend/cuda/CUDABackendFusedOps.cu b/src/backend/cuda/CUDABackendFusedOps.cu index 48c42f9..d739486 100644 --- a/src/backend/cuda/CUDABackendFusedOps.cu +++ b/src/backend/cuda/CUDABackendFusedOps.cu @@ -107,41 +107,91 @@ void CUDABackend::ComputeError(float* e, const float* z, const float* mu, size_t CHECK_CUDA_LAUNCH(); } +// One block per batch row (not one thread): SOFTMAX_BLOCK_SIZE threads +// cooperate across the nextSize class dimension via a classic +// shared-memory halving reduction, so a small batchSize no longer leaves +// almost the entire GPU idle while each of the few threads serially walks +// the whole class dimension four times. SOFTMAX_BLOCK_SIZE must stay a +// power of two (the halving reduction relies on it) and must match the +// launch configuration below exactly. +constexpr int SOFTMAX_BLOCK_SIZE = 256; + __global__ void SoftmaxCrossEntropyKernel(float* e, const float* z, const float* mu, size_t batchSize, size_t nextSize, float* rowEnergies) { - size_t b = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + size_t b = blockIdx.x; if (b >= batchSize) return; size_t base = b * nextSize; const float eps = 1e-8f; + int tid = threadIdx.x; + + __shared__ float sdata[SOFTMAX_BLOCK_SIZE]; - float max_val = mu[base]; - for (size_t j = 1; j < nextSize; ++j) - max_val = fmaxf(max_val, mu[base + j]); + // Row max, reduced across the block (numerically-stable softmax). + float local_max = -INFINITY; + for (size_t j = tid; j < nextSize; j += SOFTMAX_BLOCK_SIZE) + local_max = fmaxf(local_max, mu[base + j]); + sdata[tid] = local_max; + __syncthreads(); + for (int stride = SOFTMAX_BLOCK_SIZE / 2; stride > 0; stride >>= 1) + { + if (tid < stride) + sdata[tid] = fmaxf(sdata[tid], sdata[tid + stride]); + __syncthreads(); + } + float max_val = sdata[0]; + __syncthreads(); - float sum_exp = 0.0f; - for (size_t j = 0; j < nextSize; ++j) + // exp(mu - max) and its row sum, same reduction shape. + float local_sum = 0.0f; + for (size_t j = tid; j < nextSize; j += SOFTMAX_BLOCK_SIZE) { float ex = expf(mu[base + j] - max_val); e[base + j] = ex; - sum_exp += ex; + local_sum += ex; } - - float row_energy = 0.0f; + sdata[tid] = local_sum; + __syncthreads(); + for (int stride = SOFTMAX_BLOCK_SIZE / 2; stride > 0; stride >>= 1) + { + if (tid < stride) + sdata[tid] += sdata[tid + stride]; + __syncthreads(); + } + float sum_exp = sdata[0]; + // Every thread still needs to read sum_exp out of sdata[0] before any + // thread starts overwriting sdata[] below -- without this barrier a + // fast thread's write could stomp sdata[0] before a slow thread's read. + __syncthreads(); + + // Normalize to a probability (always), optionally reducing the + // cross-entropy row energy alongside it. bool needRowEnergy = (rowEnergies != nullptr); - for (size_t j = 0; j < nextSize; ++j) + float local_energy = 0.0f; + for (size_t j = tid; j < nextSize; j += SOFTMAX_BLOCK_SIZE) { float prob = e[base + j] / sum_exp; e[base + j] = prob; if (needRowEnergy) - row_energy -= z[base + j] * logf(prob + eps); + local_energy -= z[base + j] * logf(prob + eps); } if (needRowEnergy) - rowEnergies[b] = row_energy; + { + sdata[tid] = local_energy; + __syncthreads(); + for (int stride = SOFTMAX_BLOCK_SIZE / 2; stride > 0; stride >>= 1) + { + if (tid < stride) + sdata[tid] += sdata[tid + stride]; + __syncthreads(); + } + if (tid == 0) + rowEnergies[b] = sdata[0]; + } - for (size_t j = 0; j < nextSize; ++j) + for (size_t j = tid; j < nextSize; j += SOFTMAX_BLOCK_SIZE) e[base + j] = z[base + j] - e[base + j]; } @@ -153,9 +203,7 @@ float CUDABackend::ComputeSoftmaxCrossEntropyErrorAndEnergy(float* e, const floa if (!e || !z || !mu || !rowEnergies || batchSize == 0 || nextSize == 0) return 0.0f; - constexpr int BLOCK_SIZE = 256; - const int blocks = static_cast((batchSize + BLOCK_SIZE - 1) / BLOCK_SIZE); - SoftmaxCrossEntropyKernel<<>>( + SoftmaxCrossEntropyKernel<<<(int)batchSize, SOFTMAX_BLOCK_SIZE, 0, stream>>>( e, z, mu, batchSize, nextSize, rowEnergies); CHECK_CUDA_LAUNCH(); @@ -172,9 +220,7 @@ void CUDABackend::ComputeSoftmaxCrossEntropyError(float* e, const float* z, cons if (!e || !z || !mu || batchSize == 0 || nextSize == 0) return; - constexpr int BLOCK_SIZE = 256; - const int blocks = static_cast((batchSize + BLOCK_SIZE - 1) / BLOCK_SIZE); - SoftmaxCrossEntropyKernel<<>>( + SoftmaxCrossEntropyKernel<<<(int)batchSize, SOFTMAX_BLOCK_SIZE, 0, stream>>>( e, z, mu, batchSize, nextSize, /*rowEnergies=*/nullptr); CHECK_CUDA_LAUNCH(); } @@ -254,21 +300,17 @@ float CUDABackend::ComputePrecisionWeightedErrorAndEnergy(float* e, const float* if (!e || !z || !mu || !p || n == 0) return 0.0f; - float* energyAccum = Allocate(1); - if (!energyAccum) - return 0.0f; - cudaMemsetAsync(energyAccum, 0, sizeof(float), stream); + cudaMemsetAsync(scalarScratch, 0, sizeof(float), stream); constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); PrecisionWeightedErrorEnergyKernel<<>>(e, z, mu, p, n, width, - energyAccum); + scalarScratch); CHECK_CUDA_LAUNCH(); float total = 0.0f; - cudaMemcpyAsync(&total, energyAccum, sizeof(float), cudaMemcpyDeviceToHost, stream); + cudaMemcpyAsync(&total, scalarScratch, sizeof(float), cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); - Free(energyAccum); return total; } @@ -301,23 +343,19 @@ float CUDABackend::Sum(const float* buf, size_t n) noexcept if (!buf || n == 0) return 0.0f; - float* sumAccum = Allocate(1); - if (!sumAccum) - return 0.0f; - cudaMemsetAsync(sumAccum, 0, sizeof(float), stream); + cudaMemsetAsync(scalarScratch, 0, sizeof(float), stream); constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); - SumKernel<<>>(buf, n, sumAccum); + SumKernel<<>>(buf, n, scalarScratch); CHECK_CUDA_LAUNCH(); // Synchronous readback, same as ComputeErrorAndEnergy and friends -- // see this method's own IComputeBackend doc: never call from inside a // captured graph region. float total = 0.0f; - cudaMemcpyAsync(&total, sumAccum, sizeof(float), cudaMemcpyDeviceToHost, stream); + cudaMemcpyAsync(&total, scalarScratch, sizeof(float), cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); - Free(sumAccum); return total; } diff --git a/src/backend/cuda/CUDABackendGemm.cu b/src/backend/cuda/CUDABackendGemm.cu index 087ecd4..c49ec35 100644 --- a/src/backend/cuda/CUDABackendGemm.cu +++ b/src/backend/cuda/CUDABackendGemm.cu @@ -8,16 +8,72 @@ * CUDABackendElementwise.cu, CUDABackendActivations.cu, * CUDABackendFusedOps.cu, CUDABackendOptimizer.cu, CUDABackendConv.cu for * the rest. + * + * @note TryFusedForwardPass() dispatches RELU/LINEAR/GELU/TANH/SIGMOID + * to a fused CUTLASS GEMM, but every current caller passes LINEAR + * unconditionally and applies its real activation as a separate, + * unfused pass on the GEMM's INPUT beforehand (see + * IComputeBackend::TryFusedForwardPass's @warning). Actually fusing a + * layer's activation into this GEMM means restructuring that caller to + * feed this GEMM the pre-activation belief instead and let the epilogue + * activate the OUTPUT -- a settling-loop data-flow change, not a new + * branch here, and unverified end-to-end since this file has no way to + * be exercised by a caller yet. */ #include #ifdef DEEPITY_USE_CUDA #include #include +// Must be defined before activation.h (pulled in by +// linear_combination_sigmoid.h below) is first included anywhere in this +// translation unit: it selects Sigmoid's tanh-identity formula +// (sigmoid(x) = 0.5*tanh(x/2)+0.5) over the default 1/(1+exp(-x)). This +// engine's own CUDABackendActivations.cu sigmoidKernelInto() uses the +// same tanh-identity (and the same tanh.approx.f32 PTX instruction CUTLASS's +// fast_tanh() uses on sm_75+), so without this, GemmFwdSigmoid below would +// silently compute a numerically different (though mathematically +// equivalent) result than the unfused activation path it's meant to +// replace -- exactly the kind of fused/unfused divergence this project's +// predictive-coding settling loop can't tolerate. +#ifndef CUTLASS_USE_TANH_FOR_SIGMOID +#define CUTLASS_USE_TANH_FOR_SIGMOID +#endif +#include +#include +#include #include namespace Deep { +// GELU_taylor, not CUTLASS's erf-exact GELU: this engine's own GELU +// (Activations.h / CUDABackendActivations.cu's GeluKernelInto) is the +// tanh approximation 0.5*x*(1+tanh(sqrt(2/pi)*x*(1+0.044715*x^2))), +// bit-for-bit CUTLASS's GELU_taylor (same constants, same fast_tanh()). +// Using plain GELU here would fuse a DIFFERENT activation than the one +// every unfused call site actually computes. +using LinearCombinationGELUTaylor = cutlass::epilogue::thread::LinearCombinationGeneric< + cutlass::epilogue::thread::GELU_taylor, float, 128 / cutlass::sizeof_bits::value, float, + float, cutlass::epilogue::thread::ScaleType::Default, cutlass::FloatRoundStyle::round_to_nearest, + /*IsHeavy=*/true>; +using LinearCombinationGELUTaylorScalar = cutlass::epilogue::thread::LinearCombinationGeneric< + cutlass::epilogue::thread::GELU_taylor, float, 1, float, float, + cutlass::epilogue::thread::ScaleType::Default, cutlass::FloatRoundStyle::round_to_nearest, + /*IsHeavy=*/true>; + +// cutlass::epilogue::thread::Tanh calls fast_tanh(), the same +// tanh.approx.f32 PTX instruction (sm_75+) this engine's own tanhKernelInto() +// uses -- no dedicated linear_combination_tanh.h ships in CUTLASS, so this +// goes through the generic epilogue directly instead of a ready-made alias. +using LinearCombinationTanh = cutlass::epilogue::thread::LinearCombinationGeneric< + cutlass::epilogue::thread::Tanh, float, 128 / cutlass::sizeof_bits::value, float, float, + cutlass::epilogue::thread::ScaleType::Default, cutlass::FloatRoundStyle::round_to_nearest, + /*IsHeavy=*/true>; +using LinearCombinationTanhScalar = cutlass::epilogue::thread::LinearCombinationGeneric< + cutlass::epilogue::thread::Tanh, float, 1, float, float, + cutlass::epilogue::thread::ScaleType::Default, cutlass::FloatRoundStyle::round_to_nearest, + /*IsHeavy=*/true>; + using GemmFwdRelu = cutlass::gemm::device::Gemm< float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, @@ -52,6 +108,55 @@ using GemmFwdLinearScalar = cutlass::gemm::device::Gemm< cutlass::epilogue::thread::LinearCombination, cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; +using GemmFwdGELU = cutlass::gemm::device::Gemm< + float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, + cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, + cutlass::gemm::GemmShape<128, 128, 16>, cutlass::gemm::GemmShape<64, 64, 16>, + cutlass::gemm::GemmShape<16, 8, 8>, LinearCombinationGELUTaylor, + cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; + +using GemmFwdGELUScalar = cutlass::gemm::device::Gemm< + float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, + cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, + cutlass::gemm::GemmShape<128, 128, 16>, cutlass::gemm::GemmShape<64, 64, 16>, + cutlass::gemm::GemmShape<16, 8, 8>, LinearCombinationGELUTaylorScalar, + cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; + +using GemmFwdTanh = cutlass::gemm::device::Gemm< + float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, + cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, + cutlass::gemm::GemmShape<128, 128, 16>, cutlass::gemm::GemmShape<64, 64, 16>, + cutlass::gemm::GemmShape<16, 8, 8>, LinearCombinationTanh, + cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; + +using GemmFwdTanhScalar = cutlass::gemm::device::Gemm< + float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, + cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, + cutlass::gemm::GemmShape<128, 128, 16>, cutlass::gemm::GemmShape<64, 64, 16>, + cutlass::gemm::GemmShape<16, 8, 8>, LinearCombinationTanhScalar, + cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; + +// CUTLASS's own ready-made alias: with CUTLASS_USE_TANH_FOR_SIGMOID +// defined above, this is Sigmoid's tanh-identity branch, matching +// CUDABackendActivations.cu's sigmoidKernelInto() exactly (see the +// #define's comment up top). +using GemmFwdSigmoid = cutlass::gemm::device::Gemm< + float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, + cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, + cutlass::gemm::GemmShape<128, 128, 16>, cutlass::gemm::GemmShape<64, 64, 16>, + cutlass::gemm::GemmShape<16, 8, 8>, + cutlass::epilogue::thread::LinearCombinationSigmoid< + float, 128 / cutlass::sizeof_bits::value, float, float>, + cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; + +using GemmFwdSigmoidScalar = cutlass::gemm::device::Gemm< + float, cutlass::layout::RowMajor, float, cutlass::layout::ColumnMajor, float, + cutlass::layout::RowMajor, float, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80, + cutlass::gemm::GemmShape<128, 128, 16>, cutlass::gemm::GemmShape<64, 64, 16>, + cutlass::gemm::GemmShape<16, 8, 8>, + cutlass::epilogue::thread::LinearCombinationSigmoid, + cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 4>; + void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, float alpha, const float* A, int lda, const float* B, int ldb, float beta, float* C, int ldc) noexcept { @@ -132,6 +237,78 @@ bool CUDABackend::TryFusedForwardPass(ActivationType actType, const float* zF, c split_k_slices); return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; } + else if (actType == ActivationType::GELU) + { + if (wide) + { + GemmFwdGELU gemm_op; + GemmFwdGELU::Arguments args({batchSize, nextSize, size}, + {zF, size}, + {W, size}, + {bias, 0}, + {mu, nextSize}, + {1.0f, 1.0f}, + split_k_slices); + return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; + } + GemmFwdGELUScalar gemm_op; + GemmFwdGELUScalar::Arguments args({batchSize, nextSize, size}, + {zF, size}, + {W, size}, + {bias, 0}, + {mu, nextSize}, + {1.0f, 1.0f}, + split_k_slices); + return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; + } + else if (actType == ActivationType::TANH) + { + if (wide) + { + GemmFwdTanh gemm_op; + GemmFwdTanh::Arguments args({batchSize, nextSize, size}, + {zF, size}, + {W, size}, + {bias, 0}, + {mu, nextSize}, + {1.0f, 1.0f}, + split_k_slices); + return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; + } + GemmFwdTanhScalar gemm_op; + GemmFwdTanhScalar::Arguments args({batchSize, nextSize, size}, + {zF, size}, + {W, size}, + {bias, 0}, + {mu, nextSize}, + {1.0f, 1.0f}, + split_k_slices); + return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; + } + else if (actType == ActivationType::SIGMOID) + { + if (wide) + { + GemmFwdSigmoid gemm_op; + GemmFwdSigmoid::Arguments args({batchSize, nextSize, size}, + {zF, size}, + {W, size}, + {bias, 0}, + {mu, nextSize}, + {1.0f, 1.0f}, + split_k_slices); + return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; + } + GemmFwdSigmoidScalar gemm_op; + GemmFwdSigmoidScalar::Arguments args({batchSize, nextSize, size}, + {zF, size}, + {W, size}, + {bias, 0}, + {mu, nextSize}, + {1.0f, 1.0f}, + split_k_slices); + return gemm_op(args, nullptr, stream) == cutlass::Status::kSuccess; + } return false; } From 2c475f79b0676600c5858e47c9ba06587f251e8e Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Sun, 4 Oct 2026 12:48:02 -0400 Subject: [PATCH 12/12] Testing on gpu. Also, added gitattributes. --- .gitattributes | 39 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 39 insertions(+) create mode 100644 .gitattributes diff --git a/.gitattributes b/.gitattributes new file mode 100644 index 0000000..8576354 --- /dev/null +++ b/.gitattributes @@ -0,0 +1,39 @@ +* text=auto eol=lf + +# though text=auto should already catch these +*.cpp text eol=lf +*.cc text eol=lf +*.h text eol=lf +*.hpp text eol=lf +*.cu text eol=lf +*.cuh text eol=lf +*.py text eol=lf +*.cmake text eol=lf +CMakeLists.txt text eol=lf +*.md text eol=lf +*.txt text eol=lf +*.json text eol=lf +*.yml text eol=lf +*.yaml text eol=lf +*.toml text eol=lf +*.sh text eol=lf +.clang-format text eol=lf + +*.png binary +*.jpg binary +*.jpeg binary +*.pdf binary +*.gz binary +*.bin binary +*.so binary +*.pyd binary +*.dll binary +*.pyc binary + +# GitHub/Linguist +data/** linguist-generated=true +resources/** linguist-generated=true +**/__pycache__/** linguist-generated=true +experiments/step-tuning/** linguist-generated=true +docs/html/** linguist-generated=true +docs/latex/** linguist-generated=true \ No newline at end of file