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 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/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/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/include/deepity/backend/CPUBackend.h b/include/deepity/backend/CPUBackend.h index 1c4a57c..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; @@ -54,6 +56,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 +157,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..7757227 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; @@ -45,6 +47,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 +150,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; @@ -167,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 2d4c44c..677d58e 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 @@ -103,6 +114,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). @@ -198,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]. @@ -337,6 +367,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..2e47c10 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. @@ -165,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/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 feb0539..ec8d99f 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. /// @@ -119,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/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..1308ef4 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; @@ -105,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/include/deepity/networks/ConvPCNetwork.h b/include/deepity/networks/ConvPCNetwork.h index 49bbc57..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. @@ -107,5 +112,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..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(); @@ -233,5 +237,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..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. @@ -162,5 +167,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..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; @@ -118,5 +122,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/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/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/__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/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/CUDABackendCore.cu b/src/backend/cuda/CUDABackendCore.cu index b5dc368..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); } @@ -90,6 +93,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/backend/cuda/CUDABackendFusedOps.cu b/src/backend/cuda/CUDABackendFusedOps.cu index 4694388..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(); } @@ -208,11 +254,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 +283,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 @@ -237,21 +300,62 @@ 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, scalarScratch, sizeof(float), cudaMemcpyDeviceToHost, stream); + cudaStreamSynchronize(stream); + + 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; + + 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, 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, energyAccum, sizeof(float), cudaMemcpyDeviceToHost, stream); + cudaMemcpyAsync(&total, scalarScratch, sizeof(float), cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); - Free(energyAccum); 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; } 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 0e7e4d8..2e71230 100644 --- a/src/networks/ConvPCNetwork.cpp +++ b/src/networks/ConvPCNetwork.cpp @@ -1,5 +1,6 @@ #include #include +#include namespace Deep { @@ -80,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 @@ -116,14 +117,83 @@ 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) + { + // 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(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + 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 + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + 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 +242,80 @@ 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(false); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + backend->Synchronize(); + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/src/networks/DirectKPPCNetwork.cpp b/src/networks/DirectKPPCNetwork.cpp index f48a26e..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 { @@ -132,6 +142,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/src/networks/DiscriminativePCNetwork.cpp b/src/networks/DiscriminativePCNetwork.cpp index f449f45..8c643d3 100644 --- a/src/networks/DiscriminativePCNetwork.cpp +++ b/src/networks/DiscriminativePCNetwork.cpp @@ -2,6 +2,7 @@ #include #include #include +#include #include #include @@ -73,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() @@ -123,18 +124,80 @@ 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(false); + UpdateState(); + } + UpdatePrecision(); + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + backend->Synchronize(); + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdatePrecision(); + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdatePrecision(); + UpdateWeights(); } - UpdatePrecision(); - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -168,14 +231,66 @@ 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(false); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + backend->Synchronize(); + } + else + { + for (int t = 0; t < inferenceSteps; t++) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; t++) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/src/networks/FullConvPCNetwork.cpp b/src/networks/FullConvPCNetwork.cpp index 48adb53..a382464 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,65 @@ 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; + + // 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + backend->Synchronize(); + } + 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/FullPCNetwork.cpp b/src/networks/FullPCNetwork.cpp index ab46e37..9771947 100644 --- a/src/networks/FullPCNetwork.cpp +++ b/src/networks/FullPCNetwork.cpp @@ -164,13 +164,27 @@ 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"; } if (graphCaptured) + { backend->ReplayGraph(); + // 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 2d3a0da..e046595 100644 --- a/src/networks/GaussSeidelPCNetwork.cpp +++ b/src/networks/GaussSeidelPCNetwork.cpp @@ -1,5 +1,6 @@ #include #include +#include namespace Deep { @@ -46,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. @@ -61,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 @@ -100,11 +105,58 @@ 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(false); + 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(); + backend->Synchronize(); + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + Step(false); + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + Step(false); + 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 +166,76 @@ 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. 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(false); + 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(); + backend->Synchronize(); + } + else + { + ProjectForward(); + GetTerminalLayer()->ComputeError(false); + for (int t = 0; t < inferenceSteps; ++t) + Step(false); + UpdateWeights(); + } + } + else + { + ProjectForward(); + GetTerminalLayer()->ComputeError(false); + for (int t = 0; t < inferenceSteps; ++t) + Step(false); + 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..1f55e6e 100644 --- a/src/networks/SimpleConvPCNetwork.cpp +++ b/src/networks/SimpleConvPCNetwork.cpp @@ -1,5 +1,6 @@ #include #include +#include namespace Deep @@ -81,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 @@ -107,14 +108,70 @@ 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) + { + // 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(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + backend->Synchronize(); + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -163,17 +220,74 @@ 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(false); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + 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 + { + std::cerr << "Graph capture failed, falling back to non-graph execution for this call.\n"; + } + } + + if (graphCapturedWithProjection) + { + backend->ReplayGraph(); + backend->Synchronize(); + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); } - UpdateWeights(); + float finalEnergy = CalculateState(); GetTerminalLayer()->UnclampState(); return finalEnergy; diff --git a/src/networks/SimplePCNetwork.cpp b/src/networks/SimplePCNetwork.cpp index 3889d2f..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 { @@ -174,6 +184,7 @@ namespace Deep if (graphCaptured) { backend->ReplayGraph(); + backend->Synchronize(); } else { diff --git a/temp.py b/temp.py index 438a527..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 @@ -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/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/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/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); 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; +}