diff --git a/CMakeLists.txt b/CMakeLists.txt index 5dcdd7a..1f95933 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -144,6 +144,10 @@ endif() if(DEEPITY_ENABLE_CUDA) find_package(CUDAToolkit QUIET) if(CUDAToolkit_FOUND) + enable_language(CUDA) + set(CMAKE_CUDA_STANDARD 17) + set(CMAKE_CUDA_STANDARD_REQUIRED ON) + set(CMAKE_CUDA_ARCHITECTURES 86) message(STATUS "CUDA support enabled.") else() message(STATUS "CUDA toolkit not found. Building CPU-only.") @@ -199,10 +203,6 @@ else() set(DEEPITY_BLAS_RESOLVED FALSE) if(DEEPITY_USE_MKL) - # Modern, recommended integration: Intel's own MKLConfig.cmake, - # present if oneAPI/MKL is properly installed and sourced (e.g. - # via `source /opt/intel/oneapi/setvars.sh`). Exposes a single - # target (MKL::MKL) with both include dirs and link libraries. find_package(MKL CONFIG QUIET) if(MKL_FOUND) message(STATUS "Using Intel MKL (found via MKLConfig.cmake).") @@ -275,6 +275,10 @@ add_library(Deepity src/DirectKPPCNetwork.cpp src/ModelIO.cpp src/StreamAlignedBatcher.cpp + src/backend/Backend.cpp + src/backend/CPUBackend.cpp + src/backend/CUDABackend.cu + src/backend/Tensor.cpp ) if(DEEPITY_USE_MKL AND DEEPITY_BLAS_RESOLVED) @@ -304,6 +308,17 @@ endif() target_compile_definitions(Deepity PUBLIC SLEEF_STATIC_LIBS) +# --- Tests ----------------------------------------------------------- + +option(DEEPITY_BUILD_TESTS "Build test/verification executables" ON) + +add_executable(MatMulLargeAsymmetricVerify tests/tMatMulLargeAsymmetricVerify.cpp) +target_link_libraries(MatMulLargeAsymmetricVerify PRIVATE Deepity) +set_target_properties(MatMulLargeAsymmetricVerify PROPERTIES + RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin +) +add_test(NAME MatMulLargeAsymmetricVerify COMMAND MatMulLargeAsymmetricVerify) + # --- Compiler flags ------------------------------------------------- if(MSVC) diff --git a/bindings/pybinding.cpp b/bindings/pybinding.cpp index 275a8ba..d1b9fb2 100644 --- a/bindings/pybinding.cpp +++ b/bindings/pybinding.cpp @@ -52,8 +52,12 @@ namespace return Deep::tanh; if (act == "sigmoid") return Deep::sigmoid; + if (act == "esigmoid") + return Deep::e_sigmoid; if (act == "relu") return Deep::relu; + if (act == "gelu") + return Deep::gelu; if (act == "linear") return Deep::linear; return Deep::relu; @@ -65,8 +69,12 @@ namespace return Deep::dTanh; if (act == "dsigmoid") return Deep::dSigmoid; + if (act == "d_esigmoid") + return Deep::d_eSigmoid; if (act == "drelu") return Deep::dRelu; + if (act == "dgelu") + return Deep::dGelu; if (act == "dLinear") return Deep::dLinear; return Deep::dRelu; @@ -82,12 +90,18 @@ namespace return Deep::ActivationType::RELU; if (act == "drelu") return Deep::ActivationType::dRELU; + if (act == "gelu") + return Deep::ActivationType::GELU; + if (act == "dgelu") + return Deep::ActivationType::dGELU; if (act == "sigmoid") return Deep::ActivationType::SIGMOID; if (act == "dsigmoid") return Deep::ActivationType::dSIGMOID; if (act == "esigmoid") return Deep::ActivationType::eSIGMOID; + if (act == "d_esigmoid") + return Deep::ActivationType::d_eSIGMOID; if (act == "dlinear") return Deep::ActivationType::dLINEAR; return Deep::ActivationType::LINEAR; @@ -136,7 +150,7 @@ namespace template void BindCommonPCLayer(nb::class_ &cls, const char *className) { - cls.def("calculate_state", &LayerT::CalculateState) + cls.def("calculate_state", static_cast(&LayerT::CalculateState)) .def("update_state", &LayerT::UpdateState) .def("update_weights", &LayerT::UpdateWeights) .def("flush", &LayerT::Flush) @@ -420,30 +434,39 @@ void bind_networks(nb::module_ &m) auto simpleNetCls = nb::class_(m, "SimplePCNetwork", "Predictive Coding Network built from SimplePCLayers."); BindCommonPCNetwork(simpleNetCls, "SimplePCNetwork"); + simpleNetCls.def("__init__", [](Deep::SimplePCNetwork *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::SimplePCNetwork(batch_size, dt); }, nb::arg("batch_size"), nb::arg("device") = "cpu", "Construct a network with a fixed batch size and device (\"cpu\" or \"cuda\"/\"gpu\")."); simpleNetCls.def("add_layer", [](Deep::SimplePCNetwork &self, int size, int next_size, float lr, float ir, float lmbda, const std::string &activation, const std::string &activation_deriv) { self.AddLayer(size, next_size, lr, ir, lmbda, resolveActEnum(activation), resolveActEnum(activation_deriv)); }, nb::arg("size"), nb::arg("next_size"), nb::arg("lr") = 1e-6f, nb::arg("ir") = 0.1f, nb::arg("lmbda") = 1e-2f, nb::arg("activation") = "relu", nb::arg("activation_deriv") = "drelu", "Add a layer to the network.") .def("set_optimizer", [](Deep::SimplePCNetwork &self, const std::string &opt) { - if (opt == "ADAM") self.SetOptimizer(Deep::OptimizerType::ADAM); - else if (opt == "ADAMW") self.SetOptimizer(Deep::OptimizerType::ADAMW); - else self.SetOptimizer(Deep::OptimizerType::SGD); }, nb::arg("optimizer"), "Sets the optimizer: ADAM, ADAMW, or SGD.") - + if (opt == "ADAM") self.SetOptimizer(Deep::OptimizerType::ADAM); + else if (opt == "ADAMW") self.SetOptimizer(Deep::OptimizerType::ADAMW); + else self.SetOptimizer(Deep::OptimizerType::SGD); }, nb::arg("optimizer"), "Sets the optimizer: ADAM, ADAMW, or SGD.") .def("project_forward", &Deep::SimplePCNetwork::ProjectForward, "Seeds hidden layers from a genuine forward pass through current " "weights, instead of zero-init. Call AFTER clamp_input(), BEFORE " "the settling loop.") - - .def("train_step_with_projection", [](Deep::SimplePCNetwork &self, FloatArray x, FloatArray y, int steps) + .def("train_step_with_projection", [](Deep::SimplePCNetwork &self, FloatArray x, FloatArray y, int steps, bool computeEnergy) { - std::vector xvec(x.data(), x.data() + x.size()); - std::vector yvec(y.data(), y.data() + y.size()); - return self.TrainStepWithProjection(xvec, yvec, steps); }, nb::arg("x"), nb::arg("y"), nb::arg("steps")) +std::vector xvec(x.data(), x.data() + x.size()); +std::vector yvec(y.data(), y.data() + y.size()); +return self.TrainStepWithProjection(xvec, yvec, steps, computeEnergy); }, nb::arg("x"), nb::arg("y"), nb::arg("steps"), nb::arg("computeEnergy") = true) .def("predict_with_projection", [](Deep::SimplePCNetwork &self, FloatArray x, int steps) { - std::vector xvec(x.data(), x.data() + x.size()); - - std::vector out_beliefs = self.PredictWithProjection(xvec, steps); +std::vector xvec(x.data(), x.data() + x.size()); +std::vector out_beliefs = self.PredictWithProjection(xvec, steps); +return CopyToNewArray(out_beliefs.data(), {out_beliefs.size()}); }, nb::arg("x"), nb::arg("steps"), "Runs forward-projection init and settling loop entirely in C++, returning terminal beliefs.") + .def("randomize_weights", [](Deep::SimplePCNetwork &self, const std::string &distribution) + { + std::random_device rd; + std::mt19937 rng(rd()); + self.RandomizeWeights(rng, distribution.c_str()); }, nb::arg("distribution"), "Initialize every layer's weights using a distribution string, " + "e.g. \"normal(0, 1)\" or \"uniform(-0.3, 0.3)\"."); - return CopyToNewArray(out_beliefs.data(), {out_beliefs.size()}); }, nb::arg("x"), nb::arg("steps"), "Runs forward-projection init and settling loop entirely in C++, returning terminal beliefs."); nb::class_(m, "GaussSeidelPCNetwork", "Predictive Coding Network with Gauss-Seidel settling dynamics.") .def(nb::init(), nb::arg("batch_size")) .def("add_layer", [](Deep::GaussSeidelPCNetwork &self, int size, int next_size, float lr, float ir, float lmbda, const std::string &activation, const std::string &activation_deriv) @@ -490,7 +513,12 @@ void bind_networks(nb::module_ &m) return layers[index].get(); }, nb::rv_policy::reference_internal); nb::class_(m, "DirectKPPCNetwork", "Predictive Coding Network with Direct Kolen-Pollack feedback alignment.") - .def(nb::init(), nb::arg("batch_size")) + .def("__init__", [](Deep::DirectKPPCNetwork *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::DirectKPPCNetwork(batch_size, dt); }, nb::arg("batch_size"), nb::arg("device") = "cpu") .def("add_layer", [](Deep::DirectKPPCNetwork &self, size_t size, size_t next_size, size_t terminal_size, float lr, float ir, float fl, float lmbda, const std::string &activation, const std::string &activation_deriv) { self.AddLayer(size, next_size, terminal_size, lr, ir, fl, lmbda, resolveActEnum(activation), resolveActEnum(activation_deriv)); }, nb::arg("size"), nb::arg("next_size"), nb::arg("terminal_size"), nb::arg("lr") = 1e-6f, nb::arg("ir") = 0.1f, nb::arg("fl") = 1e-4f, nb::arg("lmbda") = 1e-2f, nb::arg("activation") = "relu", nb::arg("activation_deriv") = "drelu") .def("compile", &Deep::DirectKPPCNetwork::Compile) @@ -700,9 +728,7 @@ void bind_utilities(nb::module_ &m) { Deep::dSigmoid(x.data(), x.size()); }); } -// ============================================================================ // Main Module Entry -// ============================================================================ NB_MODULE(pydeepity, m) { m.doc() = "Deepity: A high-performance Predictive Coding library."; diff --git a/deepity_build/__pycache__/cli.cpython-312.pyc b/deepity_build/__pycache__/cli.cpython-312.pyc index 93f7c10..3d89c8d 100644 Binary files a/deepity_build/__pycache__/cli.cpython-312.pyc and b/deepity_build/__pycache__/cli.cpython-312.pyc differ diff --git a/deepity_build/__pycache__/cli.cpython-314.pyc b/deepity_build/__pycache__/cli.cpython-314.pyc index 862b8c8..7329336 100644 Binary files a/deepity_build/__pycache__/cli.cpython-314.pyc and b/deepity_build/__pycache__/cli.cpython-314.pyc differ diff --git a/deepity_build/__pycache__/cmake_runner.cpython-312.pyc b/deepity_build/__pycache__/cmake_runner.cpython-312.pyc index 4f35c5f..19af178 100644 Binary files a/deepity_build/__pycache__/cmake_runner.cpython-312.pyc and b/deepity_build/__pycache__/cmake_runner.cpython-312.pyc differ diff --git a/deepity_build/__pycache__/cmake_runner.cpython-314.pyc b/deepity_build/__pycache__/cmake_runner.cpython-314.pyc index 6775dba..b74d0a1 100644 Binary files a/deepity_build/__pycache__/cmake_runner.cpython-314.pyc and b/deepity_build/__pycache__/cmake_runner.cpython-314.pyc differ diff --git a/deepity_build/__pycache__/config.cpython-312.pyc b/deepity_build/__pycache__/config.cpython-312.pyc index 189f0fd..f56b3d2 100644 Binary files a/deepity_build/__pycache__/config.cpython-312.pyc and b/deepity_build/__pycache__/config.cpython-312.pyc differ diff --git a/deepity_build/__pycache__/config.cpython-314.pyc b/deepity_build/__pycache__/config.cpython-314.pyc index 8a06a08..90c8b2f 100644 Binary files a/deepity_build/__pycache__/config.cpython-314.pyc and b/deepity_build/__pycache__/config.cpython-314.pyc differ diff --git a/deepity_build/__pycache__/git_info.cpython-312.pyc b/deepity_build/__pycache__/git_info.cpython-312.pyc index 9a50dc3..aa35c24 100644 Binary files a/deepity_build/__pycache__/git_info.cpython-312.pyc and b/deepity_build/__pycache__/git_info.cpython-312.pyc differ diff --git a/deepity_build/__pycache__/process.cpython-312.pyc b/deepity_build/__pycache__/process.cpython-312.pyc index 2a14851..a6f0ce1 100644 Binary files a/deepity_build/__pycache__/process.cpython-312.pyc and b/deepity_build/__pycache__/process.cpython-312.pyc differ diff --git a/deepity_build/cli.py b/deepity_build/cli.py index f2917f9..b8ebb78 100644 --- a/deepity_build/cli.py +++ b/deepity_build/cli.py @@ -79,7 +79,7 @@ def _merge_pgo_profiles(config): print(f"--- PGO: profile ready: {output} ---") -def parse_args(argv: list[str] | None = None) -> argparse.Namespace: +def parse_args(argv: list[str] | None = None) -> tuple[argparse.Namespace, list[str]]: parser = argparse.ArgumentParser( description="Deepity Cross-Platform Build & Test Runner" ) @@ -190,7 +190,7 @@ def parse_args(argv: list[str] | None = None) -> argparse.Namespace: ), ) - args = parser.parse_args(argv) + args, unknown_args = parser.parse_known_args(argv) if args.list_profiles: for profile in ARCH_PROFILES.values(): @@ -206,14 +206,12 @@ def parse_args(argv: list[str] | None = None) -> argparse.Namespace: args.arch_profile = DEFAULT_ARCH_PROFILE if args.cuda is None: - # Distributed builds default to CPU-only, since a CUDA-linked binary - # isn't portable either. Everything else keeps the old default (ON). args.cuda = args.arch_profile != "distributed" - return args + return args, unknown_args -def build_config_from_args(args: argparse.Namespace) -> BuildConfig: +def build_config_from_args(args: argparse.Namespace, extra_args: list[str]) -> BuildConfig: return BuildConfig( build_type=args.build_type, jobs=args.jobs, @@ -226,6 +224,7 @@ def build_config_from_args(args: argparse.Namespace) -> BuildConfig: clean=args.clean, verbose=args.verbose, pgo=args.pgo, + extra_cmake_args=tuple(extra_args) ) @@ -349,8 +348,8 @@ def _run_pgo_workload(config: BuildConfig) -> None: def main(argv: list[str] | None = None) -> None: - args = parse_args(argv) - config = build_config_from_args(args) + args, extra_args = parse_args(argv) + config = build_config_from_args(args, extra_args) if config.clean and config.build_dir.exists(): shutil.rmtree(config.build_dir, onexc=_rmtree_onexc) diff --git a/deepity_build/cmake_runner.py b/deepity_build/cmake_runner.py index f1bdcd7..84d0c72 100644 --- a/deepity_build/cmake_runner.py +++ b/deepity_build/cmake_runner.py @@ -40,6 +40,9 @@ def configure_command(config: BuildConfig, ninja: str | None, pgo_phase: str | N if profile.msvc_flags: cmd.append(f"-DDEEPITY_MSVC_ARCH_FLAGS={profile.msvc_flags}") + if config.extra_cmake_args: + cmd.extend(config.extra_cmake_args) + if sys.platform == "win32" and ninja: # CC being unset doesn't mean clang isn't in play -- CMake can # auto-detect and pick it up on its own (as this project's own diff --git a/deepity_build/config.py b/deepity_build/config.py index c16e469..d51dca3 100644 --- a/deepity_build/config.py +++ b/deepity_build/config.py @@ -75,6 +75,7 @@ class BuildConfig: clean: bool verbose: bool pgo: bool = False + extra_cmake_args: tuple[str,...] = () build_root: Path = Path("build") @property diff --git a/deepity_build/reporting/__pycache__/__init__.cpython-312.pyc b/deepity_build/reporting/__pycache__/__init__.cpython-312.pyc index 91e0fd9..b7ffd31 100644 Binary files a/deepity_build/reporting/__pycache__/__init__.cpython-312.pyc and b/deepity_build/reporting/__pycache__/__init__.cpython-312.pyc differ diff --git a/deepity_build/reporting/__pycache__/base.cpython-312.pyc b/deepity_build/reporting/__pycache__/base.cpython-312.pyc index f610954..5e9ef2a 100644 Binary files a/deepity_build/reporting/__pycache__/base.cpython-312.pyc and b/deepity_build/reporting/__pycache__/base.cpython-312.pyc differ diff --git a/deepity_build/reporting/__pycache__/plain_reporter.cpython-312.pyc b/deepity_build/reporting/__pycache__/plain_reporter.cpython-312.pyc deleted file mode 100644 index d217868..0000000 Binary files a/deepity_build/reporting/__pycache__/plain_reporter.cpython-312.pyc and /dev/null differ diff --git a/deepity_build/reporting/__pycache__/rich_reporter.cpython-312.pyc b/deepity_build/reporting/__pycache__/rich_reporter.cpython-312.pyc index 306af05..7b9ef29 100644 Binary files a/deepity_build/reporting/__pycache__/rich_reporter.cpython-312.pyc and b/deepity_build/reporting/__pycache__/rich_reporter.cpython-312.pyc differ diff --git a/experiments/step-tuning/pydeepity/__pycache__/__init__.cpython-314.pyc b/experiments/step-tuning/pydeepity/__pycache__/__init__.cpython-314.pyc deleted file mode 100644 index d93b238..0000000 Binary files a/experiments/step-tuning/pydeepity/__pycache__/__init__.cpython-314.pyc and /dev/null differ diff --git a/imagenet.py b/imagenet.py new file mode 100644 index 0000000..c1c3f37 --- /dev/null +++ b/imagenet.py @@ -0,0 +1,208 @@ +""" +Tiny ImageNet-200 training for DKPPCN. + +v2: deeper/wider network (~63M params, up from ~15M), matching the +literature's suggestion that published, working PCN results on Tiny +ImageNet use substantially larger networks (PCX reports up to +AlexNet-scale, ~160M params) than a naive MNIST-scale architecture. +LR/FL deliberately held at the confirmed-stable 5e-5 from the previous +run, so this test isolates architecture size as the one changed +variable -- if energy stays stable and accuracy improves, that confirms +size was the real gap; if energy destabilizes again at the same LR, +that's equally informative (suggests LR needs to scale down further as +network size grows, a documented phenomenon in this literature). +""" +import numpy as np +import os +import sys +from time import perf_counter +from pydeepity import DKPPCN +from PIL import Image + +IMG_SIZE = 64 +IMG_DIM = IMG_SIZE * IMG_SIZE * 3 +DATA_DIR = "tiny-imagenet-200" + + +def load_wnids(data_dir): + with open(os.path.join(data_dir, "wnids.txt")) as f: + return [line.strip() for line in f if line.strip()] + + +def load_image_as_vector(path): + img = Image.open(path).convert("RGB") + return np.asarray(img, dtype=np.uint8).reshape(-1) + + +def load_train_set(data_dir, wnids): + print("Loading Tiny ImageNet training set (100,000 images)...") + wnid_to_idx = {w: i for i, w in enumerate(wnids)} + + X = np.zeros((len(wnids) * 500, IMG_DIM), dtype=np.uint8) + y_idx = np.zeros(len(wnids) * 500, dtype=np.int64) + + pos = 0 + for wnid in wnids: + img_dir = os.path.join(data_dir, "train", wnid, "images") + filenames = sorted(os.listdir(img_dir)) + for fname in filenames: + X[pos] = load_image_as_vector(os.path.join(img_dir, fname)) + y_idx[pos] = wnid_to_idx[wnid] + pos += 1 + print(f" loaded class {wnid} ({pos}/{len(wnids) * 500})", end="\r") + print() + + return X[:pos], y_idx[:pos], len(wnids) + + +def load_val_set(data_dir, wnids): + print("Loading Tiny ImageNet validation set (10,000 images)...") + wnid_to_idx = {w: i for i, w in enumerate(wnids)} + + annotations = {} + with open(os.path.join(data_dir, "val", "val_annotations.txt")) as f: + for line in f: + parts = line.strip().split("\t") + annotations[parts[0]] = parts[1] + + img_dir = os.path.join(data_dir, "val", "images") + filenames = sorted(annotations.keys()) + + X = np.zeros((len(filenames), IMG_DIM), dtype=np.uint8) + y_idx = np.zeros(len(filenames), dtype=np.int64) + + for i, fname in enumerate(filenames): + X[i] = load_image_as_vector(os.path.join(img_dir, fname)) + y_idx[i] = wnid_to_idx[annotations[fname]] + if i % 1000 == 0: + print(f" loaded {i}/{len(filenames)}", end="\r") + print() + + return X, y_idx + + +def to_float_batch(X_uint8_batch): + return X_uint8_batch.astype(np.float32) / 255.0 + + +def to_one_hot(y_idx_batch, n_classes, eps=0.001): + Y = np.full((len(y_idx_batch), n_classes), eps, dtype=np.float32) + Y[np.arange(len(y_idx_batch)), y_idx_batch] = 1.0 - eps + return Y + + +def main() -> None: + SEED = int(sys.argv[1]) if len(sys.argv) > 1 else 7 + EPOCHS = int(sys.argv[2]) if len(sys.argv) > 2 else 50 + INFERENCE_STEPS = int(sys.argv[3]) if len(sys.argv) > 3 else 2 + + if not os.path.isdir(DATA_DIR): + raise FileNotFoundError( + f"'{DATA_DIR}' not found in the current directory -- " + f"run this script from wherever you unzipped tiny-imagenet-200.zip." + ) + + wnids = load_wnids(DATA_DIR) + X_train_u8, y_train_idx, N_CLASSES = load_train_set(DATA_DIR, wnids) + X_val_u8, y_val_idx = load_val_set(DATA_DIR, wnids) + + BATCH_SIZE = 250 + TERMINAL_SIZE = N_CLASSES + + HIDDEN_1 = 4096 + HIDDEN_2 = 2048 + + LR = 5e-5 + IR = 0.15 + FL = 5e-5 + LMBDA = 1e-4 + DECAY_RATE = 0.94 + + print(f"\nBuilding network ({IMG_DIM}->{HIDDEN_1}->{HIDDEN_2}->{TERMINAL_SIZE}), seed={SEED}...") + net = DKPPCN(batch_size=BATCH_SIZE, device="gpu") + net.add_layer(IMG_DIM, HIDDEN_1, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="linear") + net.add_layer(HIDDEN_1, HIDDEN_2, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="sigmoid") + net.add_layer(HIDDEN_2, TERMINAL_SIZE, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="sigmoid") + net.add_layer(TERMINAL_SIZE, 0, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="linear") + net.set_optimizer("ADAM") + net.set_psi_optimizer("ADAM") + net.compile() + net.randomize_weights() + + total_params = IMG_DIM * HIDDEN_1 + HIDDEN_1 * HIDDEN_2 + HIDDEN_2 * TERMINAL_SIZE + total_params += (IMG_DIM + HIDDEN_1 + HIDDEN_2) * TERMINAL_SIZE + print(f"Approx. total parameters (W + Psi): {total_params:,}") + + print(f"\n*** TINY IMAGENET DKP-PC RUN (v2: bigger network) ***") + print(f"Training DKPPCN: {EPOCHS} epochs, inference_steps={INFERENCE_STEPS}, ") + print(f"lr={LR}, ir={IR}, fl={FL}, lmbda={LMBDA}, decay_rate={DECAY_RATE}") + print(f"NOTE: LR/FL held fixed from the previous, confirmed-stable run --") + print(f" this test isolates architecture size as the only changed variable.\n") + + rng = np.random.default_rng(SEED) + n_train = len(X_train_u8) + n_batches = n_train // BATCH_SIZE + start_time = perf_counter() + epoch_accs = [] + + for epoch in range(EPOCHS): + current_lr = LR * (DECAY_RATE ** epoch) + net.set_learning_rate(current_lr) + current_fl = FL * (DECAY_RATE ** epoch) + net.set_feedback_rate(current_fl) + + indices = rng.permutation(n_train) + epoch_energy = 0.0 + + for b in range(n_batches): + batch_idx = indices[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] + X_batch = to_float_batch(X_train_u8[batch_idx]) + Y_batch = to_one_hot(y_train_idx[batch_idx], N_CLASSES) + + energy = net.train_step(X_batch, Y_batch, INFERENCE_STEPS) + epoch_energy += energy + + N_ACC_BATCHES = 10 + correct = 0 + total = 0 + for b in range(min(N_ACC_BATCHES, len(X_val_u8) // BATCH_SIZE)): + X_batch = to_float_batch(X_val_u8[b * BATCH_SIZE:(b + 1) * BATCH_SIZE]) + y_batch = y_val_idx[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] + + preds = net.predict(X_batch, INFERENCE_STEPS).reshape(BATCH_SIZE, N_CLASSES) + pred_classes = np.argmax(preds, axis=1) + correct += np.sum(pred_classes == y_batch) + total += BATCH_SIZE + + epoch_acc = 100.0 * correct / total + epoch_accs.append(epoch_acc) + avg_energy = epoch_energy / n_batches + elapsed = perf_counter() - start_time + print(f"Epoch {epoch+1}/{EPOCHS} | Time: {elapsed:.1f}s | Acc: {epoch_acc:.2f}% | Avg energy: {avg_energy:.4f}") + + train_time = perf_counter() - start_time + print(f"\nTraining complete in {train_time:.1f}s.") + + print("\nRunning final validation evaluation (full 10,000 images)...") + correct = 0 + total = 0 + for i in range(0, len(X_val_u8), BATCH_SIZE): + X_batch_u8 = X_val_u8[i:i + BATCH_SIZE] + y_batch = y_val_idx[i:i + BATCH_SIZE] + if len(X_batch_u8) != BATCH_SIZE: + continue + + X_batch = to_float_batch(X_batch_u8) + preds = net.predict(X_batch, INFERENCE_STEPS).reshape(BATCH_SIZE, N_CLASSES) + pred_classes = np.argmax(preds, axis=1) + correct += np.sum(pred_classes == y_batch) + total += BATCH_SIZE + + val_acc = 100.0 * correct / total + print(f"\n=== Result ===") + print(f"DKPPCN Tiny ImageNet validation accuracy: {val_acc:.2f}%") + print(f"Train time: {train_time:.1f}s") + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/include/deepity/backend/Backend.h b/include/deepity/backend/Backend.h index e69de29..e72e80f 100644 --- a/include/deepity/backend/Backend.h +++ b/include/deepity/backend/Backend.h @@ -0,0 +1,10 @@ +#pragma once +#include +#include +#include + +namespace Deep +{ + /// @brief Factory function to create the appropriate backend. + std::unique_ptr CreateBackend(DeviceType device); +} \ No newline at end of file diff --git a/include/deepity/backend/CPUBackend.h b/include/deepity/backend/CPUBackend.h index e69de29..9fc0dce 100644 --- a/include/deepity/backend/CPUBackend.h +++ b/include/deepity/backend/CPUBackend.h @@ -0,0 +1,64 @@ +#pragma once +#include +#include + +namespace Deep +{ + class CPUBackend : public IComputeBackend + { + public: + CPUBackend() = default; + ~CPUBackend() override = default; + + // @remark these are no-ops for backend purposes + void BeginGraphCapture() noexcept override {} + bool EndGraphCapture() noexcept override { return true; } // nothing to fail on CPU + void ReplayGraph() noexcept override {} + + float *Allocate(size_t numFloats) override; + void Free(float *ptr) noexcept override; + void Zero(float *ptr, size_t numFloats) noexcept override; + void Copy(float *dst, const float *src, size_t numFloats) noexcept override; + void CopyFromHost(float *deviceDst, const float *hostSrc, size_t numFloats) noexcept override; + void CopyToHost(float *hostDst, const float *deviceSrc, size_t numFloats) noexcept override; + void RandomizeNormal(float *buf, size_t n, float mean, float stddev, uint32_t seed) noexcept override; + void RandomizeUniform(float *buf, size_t n, float min, float max, uint32_t seed) noexcept override; + + /// @brief Must be called once, before Compile()'s first + /// BeginGraphCapture(), for any backend that needs to prepare + /// batch-size-dependent state (e.g. CUDABackend's cached all-ones + /// vector for SumRows' GEMV). No-op on CPUBackend. + void PrepareForBatchSize(size_t batchSize) noexcept override {} + + void 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 override; + void SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept override; + + void Scale(float *buf, size_t n, float alpha) noexcept override; + void AxpyInto(float *y, const float *x, size_t n, float alpha) noexcept override; + void AddBiasBroadcast(float *buf, const float *bias, size_t batchSize, size_t width) noexcept override; + + void Activation(ActivationType type, float *buf, size_t n) noexcept override; + void ActivationInto(ActivationType type, float *dst, const float *src, size_t n) noexcept override; + void ActivationDerivative(ActivationType type, float *buf, size_t n, bool activated) noexcept override; + void ActivationDerivativeInto(ActivationType type, float *dst, const float *src, size_t n) noexcept override; + + void FusedStateUpdate(float *z, const float *feedback, const float *deriv, + const float *e, size_t n, float ir) noexcept override; + float ComputeErrorAndEnergy(float *e, const float *z, const float *mu, size_t n) noexcept override; + void ComputeError(float *e, const float *z, const float *mu, size_t n) noexcept override; + + void IncrementCounter(int *counter) noexcept override; + + void AdamStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, + float beta1 = 0.9f, float beta2 = 0.999f, float eps = 1e-8f) noexcept override; + void AdamWStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, float weightDecay, + float beta1 = 0.9f, float beta2 = 0.999f, float eps = 1e-8f) noexcept override; + + DeviceType GetDeviceType() const noexcept override { return DeviceType::DEVICE_CPU; }; + }; +} \ No newline at end of file diff --git a/include/deepity/backend/CUDABackend.h b/include/deepity/backend/CUDABackend.h new file mode 100644 index 0000000..669d028 --- /dev/null +++ b/include/deepity/backend/CUDABackend.h @@ -0,0 +1,75 @@ +#pragma once +#include +#ifdef DEEPITY_USE_CUDA +#include +#endif +#include + +namespace Deep +{ + class CUDABackend : public IComputeBackend + { + public: + CUDABackend(); + ~CUDABackend() override; + void BeginGraphCapture() noexcept override; + bool EndGraphCapture() noexcept override; + void ReplayGraph() noexcept override; + + float *Allocate(size_t numFloats) override; + void Free(float *ptr) noexcept override; + void Zero(float *ptr, size_t numFloats) noexcept override; + void Copy(float *dst, const float *src, size_t numFloats) noexcept override; + void CopyFromHost(float *deviceDst, const float *hostSrc, size_t numFloats) noexcept override; + void CopyToHost(float *hostDst, const float *deviceSrc, size_t numFloats) noexcept override; + void RandomizeNormal(float *buf, size_t n, float mean, float stddev, uint32_t seed) noexcept override; + void RandomizeUniform(float *buf, size_t n, float min, float max, uint32_t seed) noexcept override; + + void SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept override; + + /// @brief Must be called once, before Compile()'s first + /// BeginGraphCapture(), for any backend that needs to prepare + /// batch-size-dependent state (e.g. CUDABackend's cached all-ones + /// vector for SumRows' GEMV). No-op on CPUBackend. + void PrepareForBatchSize(size_t batchSize) noexcept override; + + void 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 override; + + void Scale(float *buf, size_t n, float alpha) noexcept override; + void AxpyInto(float *y, const float *x, size_t n, float alpha) noexcept override; + void AddBiasBroadcast(float *buf, const float *bias, size_t batchSize, size_t width) noexcept override; + + void Activation(ActivationType type, float *buf, size_t n) noexcept override; + void ActivationInto(ActivationType type, float *dst, const float *src, size_t n) noexcept override; + void ActivationDerivative(ActivationType type, float *buf, size_t n, bool activated) noexcept override; + void ActivationDerivativeInto(ActivationType type, float *dst, const float *src, size_t n) noexcept override; + + void FusedStateUpdate(float *z, const float *feedback, const float *deriv, + const float *e, size_t n, float ir) noexcept override; + float ComputeErrorAndEnergy(float *e, const float *z, const float *mu, size_t n) noexcept override; + void ComputeError(float *e, const float *z, const float *mu, size_t n) noexcept override; + + void IncrementCounter(int *ptr) noexcept override; + + void AdamStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, + float beta1 = 0.9f, float beta2 = 0.999f, float eps = 1e-8f) noexcept override; + void AdamWStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, float weightDecay, + float beta1 = 0.9f, float beta2 = 0.999f, float eps = 1e-8f) noexcept override; + + DeviceType GetDeviceType() const noexcept override { return DeviceType::DEVICE_GPU; } + + private: + cublasHandle_t handle; + cudaStream_t stream; + cudaGraph_t graph = nullptr; + cudaGraphExec_t graphExec = nullptr; + bool hasGraph = false; + void *workspace = nullptr; + float *onesVector = nullptr; + }; +} \ No newline at end of file diff --git a/include/deepity/backend/DeviceType.h b/include/deepity/backend/DeviceType.h new file mode 100644 index 0000000..bfc730c --- /dev/null +++ b/include/deepity/backend/DeviceType.h @@ -0,0 +1,11 @@ +#pragma once + +/// @brief This file exists to provide both Tensor.h and IComputeBackend.h with device types without introducing circular dependencies. +namespace Deep +{ + enum class DeviceType + { + DEVICE_CPU, + DEVICE_GPU + }; +} \ No newline at end of file diff --git a/include/deepity/backend/GPUBackend.h b/include/deepity/backend/GPUBackend.h deleted file mode 100644 index e69de29..0000000 diff --git a/include/deepity/backend/IComputeBackend.h b/include/deepity/backend/IComputeBackend.h index 670b451..2ebb66d 100644 --- a/include/deepity/backend/IComputeBackend.h +++ b/include/deepity/backend/IComputeBackend.h @@ -1,59 +1,47 @@ #pragma once #include #include - -/** - * @file IComputeBackend.h - * @brief Abstract interface for every core numerical operation a PC - * layer needs, so that CPUBackend and CUDABackend can be swapped behind - * one pointer -- no #ifdef DEEPITY_USE_CUDA anywhere except inside - * Backend.cpp's single factory function. - * - * Derived directly from what SimplePCLayer/DirectKPPCLayer actually call - * today (CalculateState, ComputeMuOnly, UpdateState, UpdateWeights, - * DirectFeedbackUpdate), not a generic "BLAS wrapper" -- a few operations - * below are kept as single, explicit, FUSED methods (FusedStateUpdate, - * ComputeErrorAndEnergy) specifically because decomposing them into - * separate elementwise-multiply/subtract/saxpy calls would lose the - * fusion benefit this codebase already relies on on CPU, and would cost - * even more on GPU (extra global-memory round-trips between kernel - * launches, where memory bandwidth is usually the real bottleneck). - * - * Device is fixed at Tensor/allocation time (see Tensor.h) -- nothing - * here supports moving a live buffer between CPU and GPU after creation. - * - * @warning This is a first draft, not yet implemented by either - * CPUBackend or CUDABackend. Several signatures are marked below as - * open questions -- confirm/adjust before treating this as final. - */ +#include namespace Deep { + class IComputeBackend { public: virtual ~IComputeBackend() = default; - // --- Memory ----------------------------------------------------- + // Graphs + + virtual void BeginGraphCapture() noexcept = 0; + /// @brief Ends capture and instantiates the captured graph. + /// @return true if capture and instantiation both succeeded and + /// ReplayGraph() is now safe to call; false otherwise. Callers + /// MUST check this -- silently assuming success here was the + /// cause of a real bug: a failed capture left ReplayGraph() + /// permanently doing nothing on every subsequent call, since the + /// caller had no way to know capture never actually happened. + virtual bool EndGraphCapture() noexcept = 0; + virtual void ReplayGraph() noexcept = 0; + + // Memory virtual float *Allocate(size_t numFloats) = 0; virtual void Free(float *ptr) noexcept = 0; virtual void Zero(float *ptr, size_t numFloats) noexcept = 0; - - /// @brief Device-to-device copy (both buffers already live on - /// this backend's device). virtual void Copy(float *dst, const float *src, size_t numFloats) noexcept = 0; - - /// @brief Host-to-device copy. No-op memcpy on CPUBackend; - /// cudaMemcpyHostToDevice on CUDABackend. This is the ONLY thing - /// needed for "load weights saved on CPU, run on GPU"; see the - /// constructor-time initialization pattern discussed separately. virtual void CopyFromHost(float *deviceDst, const float *hostSrc, size_t numFloats) noexcept = 0; - - /// @brief Device-to-host copy, e.g. for Save()/inspection. virtual void CopyToHost(float *hostDst, const float *deviceSrc, size_t numFloats) noexcept = 0; + virtual void RandomizeNormal(float *buf, size_t n, float mean, float stddev, uint32_t seed) noexcept = 0; + virtual void RandomizeUniform(float *buf, size_t n, float min, float max, uint32_t seed) noexcept = 0; + + /// @brief Must be called once, before Compile()'s first + /// BeginGraphCapture(), for any backend that needs to prepare + /// batch-size-dependent state (e.g. CUDABackend's cached all-ones + /// vector for SumRows' GEMV). No-op on CPUBackend. + virtual void PrepareForBatchSize(size_t batchSize) noexcept = 0; - // --- GEMM --------------------------------------------------------- + // GEMM virtual void MatMul(bool transA, bool transB, int M, int N, int K, @@ -61,56 +49,46 @@ namespace Deep const float *B, int ldb, float beta, float *C, int ldc) noexcept = 0; - // --- Elementwise scalar ops --------------------------------------- + /// @brief dst[j] = sum over b in [0,batchSize) of src[b*width + j], for + /// all j in [0,width). Replaces a batchSize-iteration loop of + /// individual AxpyInto calls -- the reduction-direction counterpart to + /// AddBiasBroadcast, still unfixed until now. On GPU this is one + /// cublasSgemv call against a cached all-ones vector, reinterpreting + /// src's row-major [batchSize,width] layout as column-major + /// [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 buf *= alpha (cblas_sscal equivalent). Used for weight - /// decay (W *= 1-lambda) today. - virtual void Scale(float *buf, size_t n, float alpha) noexcept = 0; + // Elementwise scalar ops - /// @brief y += alpha * x (cblas_saxpy equivalent). Used for bias - /// adds and Adam/AdamW's own internal accumulation today. + virtual void Scale(float *buf, size_t n, float alpha) noexcept = 0; virtual void AxpyInto(float *y, const float *x, size_t n, float alpha) noexcept = 0; + virtual void AddBiasBroadcast(float *buf, const float *bias, size_t batchSize, size_t width) noexcept = 0; - // --- Activation ----------------------------------------------- + // Activation - /// @brief In-place activation, matching Deep::relu/sigmoid/etc's - /// existing single-buffer signature. virtual void Activation(ActivationType type, float *buf, size_t n) noexcept = 0; - - /// @brief Two-buffer activation: reads src, writes phi(src) into - /// dst, src left untouched. Matches ComputeMuOnly()'s zF = phi(z) - /// pattern without needing a separate scopy first. virtual void ActivationInto(ActivationType type, float *dst, const float *src, size_t n) noexcept = 0; - - /// @brief In-place derivative, matching Deep::dRelu/dSigmoid/etc's - /// existing (buf, n, activated) signature. virtual void ActivationDerivative(ActivationType type, float *buf, size_t n, bool activated) noexcept = 0; - - /// @brief Two-buffer derivative: reads RAW src, writes f'(src) - /// into dst. Matches the dReluInto/dSigmoidInto/etc family added - /// to Activations.h this session. virtual void ActivationDerivativeInto(ActivationType type, float *dst, const float *src, size_t n) noexcept = 0; - // --- Fused PC-specific ops ------- + // Fused PC-specific ops - /// @brief z[i] += ir * (feedback[i] * deriv[i] - e[i]), for all i - /// in [0, n). Matches SimplePCLayer/DirectKPPCLayer's UpdateState() - /// fused settling-step update exactly. virtual void FusedStateUpdate(float *z, const float *feedback, const float *deriv, const float *e, size_t n, float ir) noexcept = 0; - - /// @brief e[i] = z[i] - mu[i], for all i in [0, n); returns - /// 0.5 * sum(e[i]^2). Matches CalculateState()'s error+energy - /// computation exactly (the fused AVX loop from earlier tonight). virtual float ComputeErrorAndEnergy(float *e, const float *z, const float *mu, size_t n) noexcept = 0; + virtual void ComputeError(float *e, const float *z, const float *mu, size_t n) noexcept = 0; - // --- Optimizer --------------------------------------------------- + // Optimizer + + virtual void IncrementCounter(int *counter) noexcept = 0; virtual void AdamStep(float *param, const float *grad, float *m, float *v, - size_t n, int t, float lr, + size_t n, const int *t, const float *lr, float beta1 = 0.9f, float beta2 = 0.999f, float eps = 1e-8f) noexcept = 0; virtual void AdamWStep(float *param, const float *grad, float *m, float *v, - size_t n, int t, float lr, float weightDecay, + size_t n, const int *t, const float *lr, float weightDecay, float beta1 = 0.9f, float beta2 = 0.999f, float eps = 1e-8f) noexcept = 0; + + virtual DeviceType GetDeviceType() const noexcept = 0; }; } \ No newline at end of file diff --git a/include/deepity/backend/Tensor.h b/include/deepity/backend/Tensor.h index 7fe3a5a..bfd849d 100644 --- a/include/deepity/backend/Tensor.h +++ b/include/deepity/backend/Tensor.h @@ -2,6 +2,7 @@ #include #include #include +#include /** * @file Tensor.h @@ -22,12 +23,6 @@ namespace Deep { - enum class DeviceType - { - DEVICE_CPU, - DEVICE_GPU - }; - class Tensor { public: diff --git a/include/deepity/layers/DirectKPPCLayer.h b/include/deepity/layers/DirectKPPCLayer.h index 9ce30e8..c2d78e5 100644 --- a/include/deepity/layers/DirectKPPCLayer.h +++ b/include/deepity/layers/DirectKPPCLayer.h @@ -2,12 +2,28 @@ #include #include +#include #include #include +#include #include #include #include +/** + * @file DirectKPPCLayer.h + * @brief Direct Kolen-Pollack predictive coding layer, routed through + * IComputeBackend. + * + * @note As of this revision, ComputeMuOnly()'s bias-add and + * UpdateWeights()'s bias-gradient accumulation both go through + * AddBiasBroadcast()/SumRows() instead of a per-batch-row AxpyInto loop + * -- same fix SimplePCLayer already had, ported here. biasGradScratch is + * a new, small (nextSize-length) buffer allocated unconditionally + * (regardless of optimizer) specifically for the SGD branch's + * accumulate-not-overwrite bias update, which SumRows alone can't do. + */ + namespace Deep { class DirectKPPCLayer : public Layer @@ -15,7 +31,7 @@ namespace Deep protected: size_t size; size_t nextSize; - size_t terminalSize; // The size of the final output layer (e.g., 10 for MNIST) + size_t terminalSize; size_t batchSize; float lr; @@ -28,7 +44,7 @@ namespace Deep DirectKPPCLayer *layerAbove = nullptr; DirectKPPCLayer *layerBelow = nullptr; - DirectKPPCLayer *terminalLayer = nullptr; // Direct pathway to \epsilon_L + DirectKPPCLayer *terminalLayer = nullptr; ActivationType activationType; ActivationFn activation; @@ -40,20 +56,19 @@ namespace Deep int t = 0; int tPsi = 0; - // --- Memory Pointers --- - // State + using BackendDeleter = void (*)(IComputeBackend *); + std::unique_ptr backend; + float *z = nullptr; float *e = nullptr; - // Forward Weights float *W = nullptr; float *b = nullptr; float *mu = nullptr; float *cachedMu = nullptr; - float *Psi = nullptr; // Maps \epsilon_L directly to this layer's state + float *Psi = nullptr; float *proj = nullptr; - // Forward Optimizer Buffers float *grad_W = nullptr; float *grad_b = nullptr; float *m_W = nullptr; @@ -61,53 +76,79 @@ namespace Deep float *m_b = nullptr; float *v_b = nullptr; - // Direct Feedback Optimizer Buffers float *grad_Psi = nullptr; float *m_Psi = nullptr; float *v_Psi = nullptr; - // Scratch + int *t_device = nullptr; + float *lr_device = nullptr; + int *tPsi_device = nullptr; + float *fl_device = nullptr; + float *zF = nullptr; float *zFDeriv = nullptr; float *feedbackScratch = nullptr; + /// @brief nextSize-length scratch buffer for SumRows' output in + /// UpdateWeights()'s SGD branch -- SGD needs `b += lr_batch * + /// sum(local_grad)`, an accumulate, which SumRows alone can't + /// express (it only overwrites). Allocated unconditionally + /// (regardless of which optimizer is selected) since it's cheap + /// -- at most `nextSize` floats -- and simpler than branching + /// allocation on optimizer choice for this one small buffer. + float *biasGradScratch = nullptr; + std::unique_ptr localArena; public: DirectKPPCLayer(size_t size, size_t nextSize, size_t terminalSize, size_t batchSize, float learningRate, float inferenceRate, float feedback, float lmbda, - ActivationType aType, ActivationType dType); + ActivationType aType, ActivationType dType, + IComputeBackend *backend = nullptr); ~DirectKPPCLayer() override = default; - // Setup - void BindMemory(MemoryArena &arena); + template + void BindMemory(ArenaT &arena); size_t GetRequiredFloats() const noexcept; void RandomizeWeights(std::mt19937 &seedGenerator) noexcept; - // Topology void SetLayerAbove(DirectKPPCLayer *l) noexcept { layerAbove = l; } void SetLayerBelow(DirectKPPCLayer *l) noexcept { layerBelow = l; } void SetTerminalLayer(DirectKPPCLayer *l) noexcept { terminalLayer = l; } - // Core DKP-PC Mechanics - float CalculateState() noexcept override; + /// @brief Matches Layer's virtual interface exactly (always + /// computes real energy). + float CalculateState() noexcept override { return CalculateState(true); } + /// @brief NOT a virtual override -- see SimplePCLayer's + /// identical pattern. Lets the settling loop skip the + /// cublasSdot-based energy reduction (and its capture-time + /// sync) when the caller doesn't need the value. + float CalculateState(bool needEnergy) noexcept; + void ComputeMuOnly() noexcept; - void UpdateState() noexcept override; // Will now pull from terminalLayer->GetErrors() - void UpdateWeights() noexcept override; // Must compute \Delta W AND \Delta \Psi + void UpdateState() noexcept override; + void UpdateWeights() noexcept override; void DirectFeedbackUpdate() noexcept; - // Getters / Setters void ClampState(const std::vector &inputData) noexcept; void UnclampState() noexcept; void ResetState() noexcept; + /// @brief Whether this layer is currently clamped -- needed so + /// DirectKPPCNetwork::ProjectForward() can skip overwriting an + /// already-clamped layer's z with a forward-projected guess. + /// Same real bug SimplePCNetwork::ProjectForward() had; fixed + /// here for the same reason, before it gets exercised for the + /// first time by moving ProjectForward() inside graph capture. + bool IsClamped() const noexcept { return isClamped; } + void SetOptimizer(OptimizerType o) noexcept { opt = o; } void SetPsiOptimizer(OptimizerType o) noexcept { optPsi = o; } - void SetLearningRate(float learningRate) noexcept { lr = learningRate; } + void SetLearningRate(float learningRate) noexcept; void SetInferenceRate(float inferenceRate) noexcept { ir = inferenceRate; } - void SetFeedbackRate(float feedbackRate) noexcept { fl = feedbackRate; } - void SetLambda(float lmbda) noexcept { lmbda = lmbda; } + void SetFeedbackRate(float feedbackRate) noexcept; + void SetLambda(float lmbda) noexcept { this->lmbda = lmbda; } float *GetBeliefs() noexcept override { return z; } const float *GetErrors() const noexcept override { return e; } @@ -121,4 +162,4 @@ namespace Deep size_t GetOutputSize() const noexcept override { return nextSize; } size_t GetTerminalSize() const noexcept { return terminalSize; } }; -} +} \ No newline at end of file diff --git a/include/deepity/layers/SimplePCLayer.h b/include/deepity/layers/SimplePCLayer.h index f980c8e..1c57327 100644 --- a/include/deepity/layers/SimplePCLayer.h +++ b/include/deepity/layers/SimplePCLayer.h @@ -8,7 +8,9 @@ #include #include #include +#include #include +#include /** * @file SimplePCLayer.h @@ -52,8 +54,18 @@ * every step while clamped. This is an EXACT optimization, not an * approximation -- default OFF so both behaviors coexist for direct * correctness/timing comparison before trusting it. - * @version 1.0 - * @date 2026-06-30 + * + * @note As of this revision, all math is routed through an + * IComputeBackend rather than calling cblas_/Deep::* functions + * directly, so this layer can run on either CPUBackend or CUDABackend. + * `backend` defaults to nullptr, in which case the layer constructs and + * owns its own CPUBackend internally -- every existing call site + * (SimplePCNetwork::AddLayer, nanobind bindings, hand-written C++) + * keeps working completely unchanged, silently getting today's exact + * CPU-only behavior. Passing a real backend explicitly is only needed + * for new, GPU-aware call sites. + * @version 1.1 + * @date 2026-09-05 * @author Jack Rose */ @@ -74,10 +86,17 @@ namespace Deep /// @param lmbda Weight decay (L2 regularization) coefficient /// @param act Activation function /// @param dAct Derivative of activation function + /// @param backend Compute backend to route all math through. + /// Defaults to nullptr, in which case this layer + /// constructs and owns its own CPUBackend internally -- + /// existing callers don't need to change anything. Pass a + /// real backend (owned elsewhere, e.g. by the network) to + /// run this layer on GPU. SimplePCLayer(size_t size, size_t nextSize, size_t batchSize = 1, float learningRate = 1e-6f, float inferenceRate = 0.1f, float lmbda = 1e-2f, void (*act)(float *, size_t) = relu, - void (*dAct)(float *, size_t, bool) = dRelu); + void (*dAct)(float *, size_t, bool) = dRelu, + IComputeBackend *backend = nullptr); /// @brief Constructor for a SimplePCLayer, using a named /// ActivationType instead of raw function pointers. @@ -90,9 +109,13 @@ namespace Deep /// @param lmbda Weight decay (L2 regularization) coefficient /// @param aType Activation type /// @param dType Activation derivative type + /// @param backend Compute backend to route all math through. + /// See the other constructor's doc for the default-nullptr + /// behavior. SimplePCLayer(size_t size, size_t nextSize, size_t batchSize = 1, float learningRate = 1e-6f, float inferenceRate = 0.1f, float lmbda = 1e-2f, - ActivationType aType = ActivationType::RELU, ActivationType dType = ActivationType::dRELU); + ActivationType aType = ActivationType::RELU, ActivationType dType = ActivationType::dRELU, + IComputeBackend *backend = nullptr); /// @brief Calculates the total network energy state, and this /// layer's outgoing prediction (mu). @@ -105,8 +128,11 @@ namespace Deep /// ngc-learn's documented convention exactly (was previously /// mu = phi(W@z+b), activation AFTER the transform). /// (No precision weighting -- see file-level note.) - /// @return This layer's energy contribution at the current state. - float CalculateState() noexcept override; + /// @param needEnergy Asks for energy + /// @return This layer's energy contribution at the current state, if asked for. + float CalculateState(bool needEnergy = true) noexcept; + + float CalculateState() noexcept override { return CalculateState(true); } /// @brief Computes the state derivatives for inference. /// @@ -133,6 +159,8 @@ namespace Deep /// layer's beliefs to update normally again. void UnclampState() noexcept; + bool IsClamped() const noexcept { return isClamped; } + /// @brief Returns this layer's belief buffer. /// @return Pointer to this layer's `size`-length beliefs. float *GetBeliefs() noexcept override { return z; } @@ -173,7 +201,7 @@ namespace Deep /// @brief Sets the learning rate used for weight updates. /// @param lr The new learning rate. - void SetLearningRate(float lr) noexcept { this->lr = lr; } + void SetLearningRate(float lr) noexcept; /// @brief Sets the inference rate (Euler integration step size). /// @param ir The new inference rate. void SetInferenceRate(float ir) noexcept { this->ir = ir; } @@ -233,7 +261,8 @@ namespace Deep /// @brief Randomizes this layer's weights (and biases) in place. /// @param twister The classic Mersenne Twister - void RandomizeWeights(std::mt19937 &twister) noexcept; + /// @param distribution The distribution of the randomization: (normal, uniform) + void RandomizeWeights(std::mt19937 &twister, const char *distribution = "normal") noexcept; /// @brief Returns this layer's configured activation type. ActivationType GetActivationType() const noexcept { return To_AType(activation); } @@ -247,11 +276,28 @@ namespace Deep size_t GetRequiredFloats() const noexcept; /// @brief Binds this layer's weight/state/scratch buffers into the /// supplied arena. Must be called before any other operation. - /// @param arena The MemoryArena to bind into. - void BindMemory(MemoryArena &arena); + /// Templated so either MemoryArena (CPU) or DeviceMemoryArena + /// (GPU) can be bound, resolved entirely at compile time -- see + /// the .cpp's explicit instantiations for the two concrete types + /// actually used. + /// @param arena The arena to bind into. + template + void BindMemory(ArenaT &arena); private: std::unique_ptr localArena; + + /// @brief The compute backend this layer routes all math + /// through. A unique_ptr with a swappable deleter, rather than + /// two separate owning/non-owning members: when constructed with + /// an explicit external backend (network-owned, GPU case), the + /// deleter is a no-op; when this layer had to construct its own + /// fallback CPUBackend (backend=nullptr was passed in), the + /// deleter actually frees it. Every call site still just uses + /// backend->Something(), identical to a raw pointer. + using BackendDeleter = void (*)(IComputeBackend *); + std::unique_ptr backend; + float *W; float *b; float *e; @@ -286,6 +332,9 @@ namespace Deep float lmbda; bool isClamped = false; + int *t_device = nullptr; + float *lr_device = nullptr; + float muCacheThreshold = -1.0f; // -1 = disabled. 0 = exact clamped-only // (today's validated behavior). >0 = // approximate, extends to unclamped @@ -315,4 +364,4 @@ namespace Deep friend class SimplePCNDiagnostics; }; -} // namespace Deep +} // namespace Deep \ No newline at end of file diff --git a/include/deepity/networks/DirectKPPCNetwork.h b/include/deepity/networks/DirectKPPCNetwork.h index 8fef9cd..5af466a 100644 --- a/include/deepity/networks/DirectKPPCNetwork.h +++ b/include/deepity/networks/DirectKPPCNetwork.h @@ -4,6 +4,8 @@ #include #include #include +#include +#include /** * @file DirectKPPCNetwork.h @@ -31,7 +33,13 @@ * it does not depend on any ordering among layers -- each layer's * update only reads Psi (which does not change during this phase) * and epsilon_L, both already available. Unlike GaussSeidelPCLayer, - * there is no sweep-ordering constraint here. + * there is no sweep-ordering constraint here. CONFIRMED by tracing + * every read/write in DirectFeedbackUpdate(): every layer writes + * only to its own proj/W, reads only from a shared, read-only + * terminalLayer error buffer and its own layerAbove's Psi -- zero + * cross-layer RAW/WAW hazard. This is the specific phase intended + * as the first real parallelism target (multiple CUDA streams, one + * per layer), once this class is confirmed correct on GPU. * * 2. Inference phase -- ordinary PC settling (CalculateState() + * UpdateState() across every layer), for inferenceSteps steps. @@ -47,10 +55,17 @@ * which (per DirectKPPCLayer) now updates both W (from the * settled state) and Psi (from the settled state and epsilon_L). * - * @warning Not yet gradient-checked at the network level. The - * individual layer's math should be independently verified before - * trusting any accuracy conclusion drawn from real training with this - * class. + * @note As of this revision, routed through IComputeBackend -- same + * refactor as SimplePCNetwork's own GPU port. `device` defaults to + * DEVICE_CPU, preserving existing behavior for anyone not explicitly + * requesting DEVICE_GPU. + * + * @warning Not yet gradient-checked at the network level, and not yet + * tested on GPU at all. The individual layer's math should be + * independently verified before trusting any accuracy conclusion drawn + * from real training with this class. This is Stage 1 of a three-stage + * plan (CPU-correct port -> GPU-correct, sequential -> exploit phase 1's + * confirmed cross-layer parallelism) -- this revision is Stage 1 only. */ namespace Deep @@ -60,7 +75,11 @@ namespace Deep public: /// @brief Constructs an empty network with a predetermined batch size. /// @param batchSize Batch size - explicit DirectKPPCNetwork(int batchSize) noexcept; + /// @param device Which device this network's layers should run + /// on. Defaults to DEVICE_CPU, preserving existing + /// behavior exactly for anyone not explicitly requesting + /// DEVICE_GPU. + explicit DirectKPPCNetwork(int batchSize, DeviceType device = DeviceType::DEVICE_CPU) noexcept; /// @brief Default destructor. ~DirectKPPCNetwork() = default; @@ -114,14 +133,16 @@ namespace Deep /// @brief Phase 1: runs DirectFeedbackUpdate() on every /// non-terminal layer. Order among layers does not matter (see /// class-level docs) -- unlike GaussSeidelPCNetwork, no sweep - /// ordering is required here. + /// ordering is required here. Still sequential in this + /// revision; see class-level warning. void DirectFeedbackUpdate() noexcept; /// @brief Phase 2: runs one ordinary PC settling step /// (CalculateState() + UpdateState()) across every layer. + /// @param computeEnergy asks for energy to be returned /// @return Total energy, summed from every layer's /// CalculateState(). - float Step() noexcept; + float Step(bool computeEnergy = true) noexcept; /// @brief Phase 3: updates every non-terminal layer's W and Psi. /// Called once after the settling loop completes. @@ -160,6 +181,8 @@ namespace Deep const std::vector> &GetLayers() const noexcept { return layers; } /// @brief Returns the batch size for the network's layers. int GetBatchSize() const noexcept { return batchSize; } + /// @brief Returns which device this network's layers run on. + DeviceType GetDevice() const noexcept { return device; } /// @brief Full train step: reset, clamp input+target, run all /// four DKP-PC phases in order, unclamp. @@ -178,17 +201,35 @@ namespace Deep /// perturbation and no target clamped, read the terminal's beliefs. std::vector Predict(const std::vector &x, int inferenceSteps); - /// @brief Loads all layers into one contiguous block of memory, - /// and wires layerAbove/layerBelow/terminalLayer across every - /// layer. + /// @brief Loads all layers into one contiguous block of memory + /// (MemoryArena for DEVICE_CPU, DeviceMemoryArena for + /// DEVICE_GPU), and wires layerAbove/layerBelow/terminalLayer + /// across every layer. void Compile(); private: /// @brief Every layer in the network, in the order they were added. std::vector> layers; - /// @brief The contiguous memory block backing every layer's buffers. - std::unique_ptr arena; + + /// @brief The network's own compute backend, created once at + /// construction and shared by every layer added afterward. + std::unique_ptr backend; + /// @brief Which device `backend` actually is. + DeviceType device; + + /// @brief Used when device == DEVICE_CPU. + std::unique_ptr cpuArena; + /// @brief Used when device == DEVICE_GPU. Only compiled when + /// DEEPITY_USE_CUDA is defined. +#if defined(DEEPITY_USE_CUDA) + std::unique_ptr gpuArena; +#endif + /// @brief The batch size shared by every layer in the network. int batchSize; + + // @brief CUDA Graph: + bool graphCaptured = false; + int capturedInferenceSteps = -1; }; } \ No newline at end of file diff --git a/include/deepity/networks/SimplePCNetwork.h b/include/deepity/networks/SimplePCNetwork.h index 628d6b7..63737e9 100644 --- a/include/deepity/networks/SimplePCNetwork.h +++ b/include/deepity/networks/SimplePCNetwork.h @@ -4,6 +4,8 @@ #include #include #include +#include +#include /** * @file SimplePCNetwork.h @@ -22,8 +24,15 @@ * network.CalculateState(); * * @note All layers are stored in a vector. - * @version 1.0 - * @date 2026-06-30 + * + * @note As of this revision, the network owns a single IComputeBackend + * (CPU by default, GPU if requested at construction) and passes it down + * into every layer it creates. `device` defaults to DEVICE_CPU, so every + * existing caller (nanobind bindings, hand-written C++) keeps compiling + * and behaving exactly as before -- passing DEVICE_GPU explicitly is + * only needed for new, GPU-aware call sites. + * @version 1.1 + * @date 2026-09-05 * @author Jack Rose */ @@ -35,7 +44,11 @@ namespace Deep public: /// @brief Constructs an empty network with a predetermined batch size. /// @param batchSize Batch size - explicit SimplePCNetwork(int batchSize) noexcept; + /// @param device Which device this network's layers should run + /// on. Defaults to DEVICE_CPU, preserving existing + /// behavior exactly for anyone not explicitly requesting + /// DEVICE_GPU. + explicit SimplePCNetwork(int batchSize, DeviceType device = DeviceType::DEVICE_CPU) noexcept; SimplePCNetwork(const SimplePCNetwork &) = delete; SimplePCNetwork &operator=(const SimplePCNetwork &) = delete; @@ -56,6 +69,11 @@ namespace Deep void AddLayer(int size, int nextSize, float lr, float ir, float lmbda, ActivationType aType, ActivationType dType); + /// @brief Randomizes the weights of each layer + /// @param rng The classic Mersenne Twister + /// @param distribution A string representation of (normal, uniform) distribution. For example: "normal(0, 1)" for a standard normal. + void RandomizeWeights(std::mt19937 &rng, const char *distribution); + /// @brief Randomizes the weights of each layer /// @param rng The classic Mersenne Twister void RandomizeWeights(std::mt19937 &rng); @@ -68,8 +86,9 @@ namespace Deep void Clamp(const std::vector &input); /// @brief Calculates the state of each layer - /// @return Returns total energy - float CalculateState(); + /// @param needEnergy ask for energy after + /// @return Returns total energy if asked for + float CalculateState(bool needEnergy = true); /// @brief Updates each layer's state void UpdateState(); @@ -95,6 +114,9 @@ namespace Deep /// @return size_t batchSize int GetBatchSize() const noexcept { return batchSize; } + /// @brief Returns which device this network's layers run on. + DeviceType GetDevice() const noexcept { return device; } + /// @brief Sets the optimizer used for weight updates on every layer. /// @param o The optimizer type to apply. void SetOptimizer(OptimizerType o) noexcept @@ -146,20 +168,47 @@ namespace Deep /// calls, update_weights, unclamp_state -- over 40 individual /// crossings per batch at STEPS=20). This does the whole sequence in /// ONE crossing instead. - float TrainStepWithProjection(const std::vector &x, const std::vector &y, int inferenceSteps); + float TrainStepWithProjection(const std::vector &x, const std::vector &y, + int inferenceSteps, bool computeEnergy = true); - std::vector PredictWithProjection(const std::vector &x, int inferenceSteps); + std::vector PredictWithProjection(const std::vector &x, int inferenceSteps); /// @brief Sets mu-cache threshold on every layer -- see /// SimplePCLayer::SetMuCacheThreshold() for semantics. Safe to call any /// time after Compile(). void SetMuCacheThreshold(float threshold) noexcept; - /// @brief Loads all layers into one contiguous block of memory. + /// @brief Loads all layers into one contiguous block of memory + /// (MemoryArena for DEVICE_CPU, DeviceMemoryArena for DEVICE_GPU). void Compile(); private: std::vector> layers; - std::unique_ptr arena; + + /// @brief The network's own compute backend, created once at + /// construction and shared by every layer added afterward. + std::unique_ptr backend; + /// @brief Which device `backend` actually is -- kept alongside + /// it since IComputeBackend itself doesn't expose its own type, + /// and Compile() needs to know which arena type to construct. + DeviceType device; + + bool graphCaptured = false; + int capturedInferenceSteps = -1; + + /// @brief Used when device == DEVICE_CPU. Only one of + /// cpuArena/gpuArena is ever actually constructed for a given + /// network -- they're kept as two separate members (rather than + /// one unified type) because MemoryArena and DeviceMemoryArena + /// share no common base, matching the same reasoning already + /// applied when DeviceMemoryArena was designed. + std::unique_ptr cpuArena; + /// @brief Used when device == DEVICE_GPU. Only compiled at all + /// when DEEPITY_USE_CUDA is defined, matching + /// DeviceMemoryArena.h's own guard. +#if defined(DEEPITY_USE_CUDA) + std::unique_ptr gpuArena; +#endif + int batchSize; }; -} +} \ No newline at end of file diff --git a/include/deepity/utils/Activations.h b/include/deepity/utils/Activations.h index 0254b90..1b75a87 100644 --- a/include/deepity/utils/Activations.h +++ b/include/deepity/utils/Activations.h @@ -50,9 +50,12 @@ namespace Deep { RELU, dRELU, + GELU, + dGELU, SIGMOID, dSIGMOID, eSIGMOID, + d_eSIGMOID, TANH, dTANH, LINEAR, @@ -61,18 +64,23 @@ namespace Deep }; static inline void relu(float *, size_t) noexcept; + static inline void gelu(float *, size_t) noexcept; static inline void sigmoid(float *, size_t) noexcept; static inline void e_sigmoid(float *, size_t) noexcept; static inline void tanh(float *, size_t) noexcept; static inline void linear(float *, size_t) noexcept; static inline void dRelu(float *, size_t, bool) noexcept; + static inline void dGelu(float *, size_t, bool) noexcept; static inline void dSigmoid(float *, size_t, bool) noexcept; + static inline void d_eSigmoid(float *, size_t, bool) noexcept; static inline void dTanh(float *, size_t, bool) noexcept; static inline void dLinear(float *, size_t, bool) noexcept; static inline void dReluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; static inline void dSigmoidInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + static inline void d_eSigmoidInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; static inline void dTanhInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; static inline void dLinearInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; @@ -82,6 +90,8 @@ namespace Deep { case ActivationType::RELU: return relu; + case ActivationType::GELU: + return gelu; case ActivationType::SIGMOID: return sigmoid; case ActivationType::eSIGMOID: @@ -102,8 +112,12 @@ namespace Deep { case ActivationType::dRELU: return dRelu; + case ActivationType::dGELU: + return dGelu; case ActivationType::dSIGMOID: return dSigmoid; + case ActivationType::d_eSIGMOID: + return d_eSigmoid; case ActivationType::dTANH: return dTanh; case ActivationType::dLINEAR: @@ -122,8 +136,12 @@ namespace Deep { case ActivationType::dRELU: return dReluInto; + case ActivationType::dGELU: + return dGeluInto; case ActivationType::dSIGMOID: return dSigmoidInto; + case ActivationType::d_eSIGMOID: + return d_eSigmoidInto; case ActivationType::dTANH: return dTanhInto; case ActivationType::dLINEAR: @@ -138,6 +156,8 @@ namespace Deep { if (fn == relu) return ActivationType::RELU; + if (fn == gelu) + return ActivationType::GELU; if (fn == sigmoid) return ActivationType::SIGMOID; if (fn == e_sigmoid) @@ -153,8 +173,12 @@ namespace Deep { if (dfn == dRelu) return ActivationType::dRELU; + if (dfn == dGelu) + return ActivationType::dGELU; if (dfn == dSigmoid) return ActivationType::dSIGMOID; + if (dfn == d_eSigmoid) + return ActivationType::d_eSIGMOID; if (dfn == dTanh) return ActivationType::dTANH; if (dfn == dLinear) @@ -509,19 +533,319 @@ namespace Deep } #pragma endregion -#define TANH_TINYLIMIT 0.000244140625f -#define TANH_POLY_LIMIT 0.625f -#define TANH_BIGLIMIT 6.0f +#pragma region gelu -#define P0 -0.9643991794f -#define P1 -99.28772310f -#define P2 -1614.687684f + constexpr float MAGIC_GELU_1 = 0.7978845608028654f; + constexpr float MAGIC_GELU_2 = 0.044715f; -#define Q0 112.8116785f -#define Q1 2235.488391f -#define Q2 4844.063053f + static inline void gelu(float *RESTRICT x, const size_t n) noexcept + { + assert(n != 0 && "n must not be 0."); + assert(x != nullptr && "x must not be null."); + + ptrdiff_t simd_end = 0; + +#if defined(__AVX512F__) + const __m512 ones512 = _mm512_set1_ps(1.0f); + const __m512 half512 = _mm512_set1_ps(0.5f); + const __m512 sqrt2overpi512 = _mm512_set1_ps(MAGIC_GELU_1); // C1 = sqrt(2/pi) + const __m512 gelu_coeff512 = _mm512_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); // C2 = C1 * 0.044715 + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)15; + for (ptrdiff_t i = 0; i < simd_end; i += 16) + { + __m512 x512 = _mm512_load_ps(x + i); + __m512 cube512 = _mm512_mul_ps(x512, _mm512_mul_ps(x512, x512)); + __m512 c1x512 = _mm512_mul_ps(sqrt2overpi512, x512); // C1*x + __m512 inner512 = _mm512_fmadd_ps(cube512, gelu_coeff512, c1x512); // x^3*C2 + C1*x + __m512 t512 = Sleef_tanhf16_u10(inner512); + __m512 res = _mm512_mul_ps(half512, _mm512_mul_ps(x512, _mm512_add_ps(t512, ones512))); + _mm512_store_ps(x + i, res); + } +#elif defined(__AVX__) + const __m256 ones256 = _mm256_set1_ps(1.0f); + const __m256 half256 = _mm256_set1_ps(0.5f); + const __m256 sqrt2overpi256 = _mm256_set1_ps(MAGIC_GELU_1); // C1 = sqrt(2/pi) + const __m256 gelu_coeff256 = _mm256_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); // C2 = C1 * 0.044715 + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)7; + for (ptrdiff_t i = 0; i < simd_end; i += 8) + { + __m256 x256 = _mm256_load_ps(x + i); + __m256 cube256 = _mm256_mul_ps(x256, _mm256_mul_ps(x256, x256)); + __m256 c1x256 = _mm256_mul_ps(sqrt2overpi256, x256); // C1*x +#if defined(__AVX2__) + __m256 inner256 = _mm256_fmadd_ps(cube256, gelu_coeff256, c1x256); // x^3*C2 + C1*x +#else + __m256 inner256 = _mm256_add_ps(_mm256_mul_ps(cube256, gelu_coeff256), c1x256); +#endif + __m256 t256 = Sleef_tanhf8_u10(inner256); + __m256 res = _mm256_mul_ps(half256, _mm256_mul_ps(x256, _mm256_add_ps(t256, ones256))); + _mm256_store_ps(x + i, res); + } +#elif defined(__SSE__) || defined(_M_AMD64) || defined(_M_X64) + const __m128 ones128 = _mm_set1_ps(1.0f); + const __m128 half128 = _mm_set1_ps(0.5f); + const __m128 sqrt2overpi128 = _mm_set1_ps(MAGIC_GELU_1); // C1 = sqrt(2/pi) + const __m128 gelu_coeff128 = _mm_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); // C2 = C1 * 0.044715 + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)3; + for (ptrdiff_t i = 0; i < simd_end; i += 4) + { + __m128 x128 = _mm_load_ps(x + i); + __m128 cube128 = _mm_mul_ps(x128, _mm_mul_ps(x128, x128)); + __m128 c1x128 = _mm_mul_ps(sqrt2overpi128, x128); // C1*x +#ifdef __FMA__ + __m128 inner128 = _mm_fmadd_ps(cube128, gelu_coeff128, c1x128); // x^3*C2 + C1*x +#else + __m128 inner128 = _mm_add_ps(c1x128, _mm_mul_ps(cube128, gelu_coeff128)); +#endif + __m128 t128 = Sleef_tanhf4_u10(inner128); + __m128 res = _mm_mul_ps(half128, _mm_mul_ps(x128, _mm_add_ps(t128, ones128))); + _mm_store_ps(x + i, res); + } +#endif + for (ptrdiff_t i = simd_end; i < n; ++i) + { + float xi = x[i]; + float inner = MAGIC_GELU_1 * xi + (MAGIC_GELU_2 * MAGIC_GELU_1) * xi * xi * xi; + x[i] = 0.5f * xi * (1.0f + Sleef_tanhf_u10(inner)); + } + } + static inline void dGelu(float *RESTRICT x, const size_t n, const bool activated) noexcept + { + assert(n != 0 && "n must not be 0."); + assert(x != nullptr && "x must not be null."); + + ptrdiff_t simd_end = 0; + +#if defined(__AVX512F__) + const __m512 ones512 = _mm512_set1_ps(1.0f); + const __m512 half512 = _mm512_set1_ps(0.5f); + const __m512 c1_512 = _mm512_set1_ps(MAGIC_GELU_1); // C1 + const __m512 c2_512 = _mm512_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); // C2 + const __m512 c2_x3_512 = _mm512_set1_ps(3.0f * MAGIC_GELU_2 * MAGIC_GELU_1); // 3 * C2 + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)15; + for (ptrdiff_t i = 0; i < simd_end; i += 16) + { + __m512 x512 = _mm512_load_ps(x + i); + __m512 xsq512 = _mm512_mul_ps(x512, x512); + __m512 xcube512 = _mm512_mul_ps(xsq512, x512); + __m512 c1x512 = _mm512_mul_ps(c1_512, x512); + __m512 inner512 = _mm512_fmadd_ps(xcube512, c2_512, c1x512); + + __m512 t512 = Sleef_tanhf16_u10(inner512); + __m512 tsq512 = _mm512_mul_ps(t512, t512); + + __m512 one_minus_tsq512 = _mm512_sub_ps(ones512, tsq512); + __m512 gprime512 = _mm512_fmadd_ps(xsq512, c2_x3_512, c1_512); + + __m512 term1 = _mm512_mul_ps(half512, _mm512_add_ps(ones512, t512)); + __m512 term2 = _mm512_mul_ps(half512, _mm512_mul_ps(x512, _mm512_mul_ps(gprime512, one_minus_tsq512))); + __m512 res = _mm512_add_ps(term1, term2); + _mm512_store_ps(x + i, res); + } +#elif defined(__AVX__) + const __m256 ones256 = _mm256_set1_ps(1.0f); + const __m256 half256 = _mm256_set1_ps(0.5f); + const __m256 c1_256 = _mm256_set1_ps(MAGIC_GELU_1); + const __m256 c2_256 = _mm256_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); + const __m256 c2_x3_256 = _mm256_set1_ps(3.0f * MAGIC_GELU_2 * MAGIC_GELU_1); + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)7; + for (ptrdiff_t i = 0; i < simd_end; i += 8) + { + __m256 x256 = _mm256_load_ps(x + i); + __m256 xsq256 = _mm256_mul_ps(x256, x256); + __m256 xcube256 = _mm256_mul_ps(xsq256, x256); + __m256 c1x256 = _mm256_mul_ps(c1_256, x256); + +#if defined(__AVX2__) + __m256 inner256 = _mm256_fmadd_ps(xcube256, c2_256, c1x256); + __m256 gprime256 = _mm256_fmadd_ps(xsq256, c2_x3_256, c1_256); +#else + __m256 inner256 = _mm256_add_ps(_mm256_mul_ps(xcube256, c2_256), c1x256); + __m256 gprime256 = _mm256_add_ps(_mm256_mul_ps(xsq256, c2_x3_256), c1_256); +#endif + + __m256 t256 = Sleef_tanhf8_u10(inner256); + __m256 tsq256 = _mm256_mul_ps(t256, t256); + __m256 one_minus_tsq256 = _mm256_sub_ps(ones256, tsq256); + + __m256 term1 = _mm256_mul_ps(half256, _mm256_add_ps(ones256, t256)); + __m256 term2 = _mm256_mul_ps(half256, _mm256_mul_ps(x256, _mm256_mul_ps(gprime256, one_minus_tsq256))); + + _mm256_store_ps(x + i, _mm256_add_ps(term1, term2)); + } +#elif defined(__SSE__) || defined(_M_AMD64) || defined(_M_X64) + const __m128 ones128 = _mm_set1_ps(1.0f); + const __m128 half128 = _mm_set1_ps(0.5f); + const __m128 c1_128 = _mm_set1_ps(MAGIC_GELU_1); + const __m128 c2_128 = _mm_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); + const __m128 c2_x3_128 = _mm_set1_ps(3.0f * MAGIC_GELU_2 * MAGIC_GELU_1); + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)3; + for (ptrdiff_t i = 0; i < simd_end; i += 4) + { + __m128 x128 = _mm_load_ps(x + i); + __m128 xsq128 = _mm_mul_ps(x128, x128); + __m128 xcube128 = _mm_mul_ps(xsq128, x128); + __m128 c1x128 = _mm_mul_ps(c1_128, x128); + +#ifdef __FMA__ + __m128 inner128 = _mm_fmadd_ps(xcube128, c2_128, c1x128); + __m128 gprime128 = _mm_fmadd_ps(xsq128, c2_x3_128, c1_128); +#else + __m128 inner128 = _mm_add_ps(_mm_mul_ps(xcube128, c2_128), c1x128); + __m128 gprime128 = _mm_add_ps(_mm_mul_ps(xsq128, c2_x3_128), c1_128); +#endif + + __m128 t128 = Sleef_tanhf4_u10(inner128); + __m128 tsq128 = _mm_mul_ps(t128, t128); + __m128 one_minus_tsq128 = _mm_sub_ps(ones128, tsq128); + + __m128 term1 = _mm_mul_ps(half128, _mm_add_ps(ones128, t128)); + __m128 term2 = _mm_mul_ps(half128, _mm_mul_ps(x128, _mm_mul_ps(gprime128, one_minus_tsq128))); + + _mm_store_ps(x + i, _mm_add_ps(term1, term2)); + } +#endif + for (ptrdiff_t i = simd_end; i < n; ++i) + { + float xi = x[i]; + float xsq = xi * xi; + + float inner = MAGIC_GELU_1 * xi + (MAGIC_GELU_2 * MAGIC_GELU_1) * xsq * xi; + float t = Sleef_tanhf_u10(inner); + + float gprime = MAGIC_GELU_1 + (3.0f * MAGIC_GELU_2 * MAGIC_GELU_1) * xsq; + + float term1 = 0.5f * (1.0f + t); + float term2 = 0.5f * xi * gprime * (1.0f - t * t); + + x[i] = term1 + term2; + } + } + + static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + { + + ptrdiff_t simd_end = 0; + +#if defined(__AVX512F__) + const __m512 ones512 = _mm512_set1_ps(1.0f); + const __m512 half512 = _mm512_set1_ps(0.5f); + const __m512 c1_512 = _mm512_set1_ps(MAGIC_GELU_1); // C1 + const __m512 c2_512 = _mm512_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); // C2 + const __m512 c2_x3_512 = _mm512_set1_ps(3.0f * MAGIC_GELU_2 * MAGIC_GELU_1); // 3 * C2 + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)15; + for (ptrdiff_t i = 0; i < simd_end; i += 16) + { + __m512 x512 = _mm512_load_ps(src + i); + __m512 xsq512 = _mm512_mul_ps(x512, x512); + __m512 xcube512 = _mm512_mul_ps(xsq512, x512); + __m512 c1x512 = _mm512_mul_ps(c1_512, x512); + __m512 inner512 = _mm512_fmadd_ps(xcube512, c2_512, c1x512); + + __m512 t512 = Sleef_tanhf16_u10(inner512); + __m512 tsq512 = _mm512_mul_ps(t512, t512); + + __m512 one_minus_tsq512 = _mm512_sub_ps(ones512, tsq512); + __m512 gprime512 = _mm512_fmadd_ps(xsq512, c2_x3_512, c1_512); + + __m512 term1 = _mm512_mul_ps(half512, _mm512_add_ps(ones512, t512)); + __m512 term2 = _mm512_mul_ps(half512, _mm512_mul_ps(x512, _mm512_mul_ps(gprime512, one_minus_tsq512))); + __m512 res = _mm512_add_ps(term1, term2); + _mm512_store_ps(dst + i, res); + } +#elif defined(__AVX__) + const __m256 ones256 = _mm256_set1_ps(1.0f); + const __m256 half256 = _mm256_set1_ps(0.5f); + const __m256 c1_256 = _mm256_set1_ps(MAGIC_GELU_1); + const __m256 c2_256 = _mm256_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); + const __m256 c2_x3_256 = _mm256_set1_ps(3.0f * MAGIC_GELU_2 * MAGIC_GELU_1); + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)7; + for (ptrdiff_t i = 0; i < simd_end; i += 8) + { + __m256 x256 = _mm256_load_ps(src + i); + __m256 xsq256 = _mm256_mul_ps(x256, x256); + __m256 xcube256 = _mm256_mul_ps(xsq256, x256); + __m256 c1x256 = _mm256_mul_ps(c1_256, x256); + +#if defined(__AVX2__) + __m256 inner256 = _mm256_fmadd_ps(xcube256, c2_256, c1x256); + __m256 gprime256 = _mm256_fmadd_ps(xsq256, c2_x3_256, c1_256); +#else + __m256 inner256 = _mm256_add_ps(_mm256_mul_ps(xcube256, c2_256), c1x256); + __m256 gprime256 = _mm256_add_ps(_mm256_mul_ps(xsq256, c2_x3_256), c1_256); +#endif + + __m256 t256 = Sleef_tanhf8_u10(inner256); + __m256 tsq256 = _mm256_mul_ps(t256, t256); + __m256 one_minus_tsq256 = _mm256_sub_ps(ones256, tsq256); + + __m256 term1 = _mm256_mul_ps(half256, _mm256_add_ps(ones256, t256)); + __m256 term2 = _mm256_mul_ps(half256, _mm256_mul_ps(x256, _mm256_mul_ps(gprime256, one_minus_tsq256))); + + _mm256_store_ps(dst + i, _mm256_add_ps(term1, term2)); + } +#elif defined(__SSE__) || defined(_M_AMD64) || defined(_M_X64) + const __m128 ones128 = _mm_set1_ps(1.0f); + const __m128 half128 = _mm_set1_ps(0.5f); + const __m128 c1_128 = _mm_set1_ps(MAGIC_GELU_1); + const __m128 c2_128 = _mm_set1_ps(MAGIC_GELU_2 * MAGIC_GELU_1); + const __m128 c2_x3_128 = _mm_set1_ps(3.0f * MAGIC_GELU_2 * MAGIC_GELU_1); + + simd_end = (ptrdiff_t)n & ~(ptrdiff_t)3; + for (ptrdiff_t i = 0; i < simd_end; i += 4) + { + __m128 x128 = _mm_load_ps(src + i); + __m128 xsq128 = _mm_mul_ps(x128, x128); + __m128 xcube128 = _mm_mul_ps(xsq128, x128); + __m128 c1x128 = _mm_mul_ps(c1_128, x128); + +#ifdef __FMA__ + __m128 inner128 = _mm_fmadd_ps(xcube128, c2_128, c1x128); + __m128 gprime128 = _mm_fmadd_ps(xsq128, c2_x3_128, c1_128); +#else + __m128 inner128 = _mm_add_ps(_mm_mul_ps(xcube128, c2_128), c1x128); + __m128 gprime128 = _mm_add_ps(_mm_mul_ps(xsq128, c2_x3_128), c1_128); +#endif + + __m128 t128 = Sleef_tanhf4_u10(inner128); + __m128 tsq128 = _mm_mul_ps(t128, t128); + __m128 one_minus_tsq128 = _mm_sub_ps(ones128, tsq128); + + __m128 term1 = _mm_mul_ps(half128, _mm_add_ps(ones128, t128)); + __m128 term2 = _mm_mul_ps(half128, _mm_mul_ps(x128, _mm_mul_ps(gprime128, one_minus_tsq128))); + + _mm_store_ps(dst + i, _mm_add_ps(term1, term2)); + } +#endif + for (ptrdiff_t i = simd_end; i < n; ++i) + { + float xi = src[i]; + float xsq = xi * xi; + + float inner = MAGIC_GELU_1 * xi + (MAGIC_GELU_2 * MAGIC_GELU_1) * xsq * xi; + float t = Sleef_tanhf_u10(inner); + + float gprime = MAGIC_GELU_1 + (3.0f * MAGIC_GELU_2 * MAGIC_GELU_1) * xsq; + + float term1 = 0.5f * (1.0f + t); + float term2 = 0.5f * xi * gprime * (1.0f - t * t); + + dst[i] = term1 + term2; + } + } + +#pragma endregion #pragma region tanh + static inline void tanh(float *RESTRICT x, const size_t n) noexcept { assert(n != 0 && "n must not be 0."); @@ -530,106 +854,279 @@ namespace Deep size_t simd_end = 0; #if defined(__AVX512F__) - simd_end = n - (n % 16); -#pragma omp parallel for schedule(static) if (n > 65536 && !omp_in_parallel()) - for (ptrdiff_t i = 0; i < (ptrdiff_t)(simd_end); i += 16) + + simd_end = n & ~(size_t)15; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t i = 0; i < static_cast(simd_end); i += 16) { - __m512 x_512 = _mm512_loadu_ps(x + i); - __m512 res = Sleef_tanhf16_u10avx512f(x_512); - _mm512_storeu_ps(x + i, res); + __m512 v = _mm512_loadu_ps(x + i); + v = Sleef_tanhf16_u10avx512f(v); + _mm512_storeu_ps(x + i, v); } #elif defined(__AVX2__) - simd_end = n - (n % 8); -#pragma omp parallel for schedule(static) if (n > 65536 && !omp_in_parallel()) - for (ptrdiff_t i = 0; i < (ptrdiff_t)(simd_end); i += 8) + + simd_end = n & ~(size_t)7; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t i = 0; i < static_cast(simd_end); i += 8) { - __m256 x_256 = _mm256_loadu_ps(x + i); - __m256 res = Sleef_tanhf8_u10avx2(x_256); - _mm256_storeu_ps(x + i, res); + __m256 v = _mm256_loadu_ps(x + i); + v = Sleef_tanhf8_u10avx2(v); + _mm256_storeu_ps(x + i, v); + } + +#elif defined(__AVX__) + + simd_end = n & ~(size_t)7; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t i = 0; i < static_cast(simd_end); i += 8) + { + __m256 v = _mm256_loadu_ps(x + i); + v = Sleef_tanhf8_u10avx(v); + _mm256_storeu_ps(x + i, v); } #elif defined(__SSE4_1__) || defined(__SSE2__) || defined(_M_AMD64) || defined(_M_X64) - simd_end = n - (n % 4); -#pragma omp parallel for schedule(static) if (n > 65536 && !omp_in_parallel()) - for (ptrdiff_t i = 0; i < (ptrdiff_t)(simd_end); i += 4) + + simd_end = n & ~(size_t)3; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t i = 0; i < static_cast(simd_end); i += 4) { - __m128 x_128 = _mm_loadu_ps(x + i); + __m128 v = _mm_loadu_ps(x + i); + #if defined(__SSE4_1__) - __m128 res = Sleef_tanhf4_u10sse4(x_128); + v = Sleef_tanhf4_u10sse4(v); #else - __m128 res = Sleef_tanhf4_u10sse2(x_128); + v = Sleef_tanhf4_u10sse2(v); #endif - _mm_storeu_ps(x + i, res); + + _mm_storeu_ps(x + i, v); } + #endif - for (size_t i = simd_end; i < n; i++) + for (size_t i = simd_end; i < n; ++i) { x[i] = Sleef_tanhf_u10(x[i]); } } - - static inline void dTanh(float *RESTRICT x, const size_t n, bool activated = false) noexcept + static inline void dTanh(float *RESTRICT x, + const size_t n, + const bool activated = false) noexcept { assert(n != 0 && "n must not be 0."); assert(x != nullptr && "x must not be null."); - if (!activated) - tanh(x, n); - size_t i = 0; - [[maybe_unused]] size_t simd_end; + + if (activated) + { +#if defined(__AVX512F__) + + const __m512 ones = _mm512_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)15; + + for (; i < simd_end; i += 16) + { + const __m512 t = _mm512_loadu_ps(x + i); + +#ifdef __FMA__ + const __m512 res = _mm512_fnmadd_ps(t, t, ones); +#else + const __m512 res = _mm512_sub_ps( + ones, + _mm512_mul_ps(t, t)); +#endif + + _mm512_storeu_ps(x + i, res); + } + +#elif defined(__AVX2__) + + const __m256 ones = _mm256_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)7; + + for (; i < simd_end; i += 8) + { + const __m256 t = _mm256_loadu_ps(x + i); + +#ifdef __FMA__ + const __m256 res = _mm256_fnmadd_ps(t, t, ones); +#else + const __m256 res = _mm256_sub_ps( + ones, + _mm256_mul_ps(t, t)); +#endif + + _mm256_storeu_ps(x + i, res); + } + +#elif defined(__SSE2__) || defined(_M_AMD64) || defined(_M_X64) + + const __m128 ones = _mm_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)3; + + for (; i < simd_end; i += 4) + { + const __m128 t = _mm_loadu_ps(x + i); + +#ifdef __FMA__ + const __m128 res = _mm_fnmadd_ps(t, t, ones); +#else + const __m128 res = _mm_sub_ps( + ones, + _mm_mul_ps(t, t)); +#endif + + _mm_storeu_ps(x + i, res); + } + +#endif + + for (; i < n; ++i) + { + const float t = x[i]; + x[i] = 1.0f - t * t; + } + + return; + } #if defined(__AVX512F__) - __m512 ones = _mm512_set1_ps(1.0f); - simd_end = n - (n % 16); - for (; i < simd_end; i += 16) { - __m512 t = _mm512_loadu_ps(x + i); + const __m512 ones = _mm512_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)15; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; + j < static_cast(simd_end); + j += 16) + { + __m512 t = _mm512_loadu_ps(x + j); + t = Sleef_tanhf16_u10avx512f(t); #ifdef __FMA__ - __m512 res = _mm512_fnmadd_ps(t, t, ones); // 1 - t*t + t = _mm512_fnmadd_ps(t, t, ones); #else - __m512 res = _mm512_sub_ps(ones, _mm512_mul_ps(t, t)); + t = _mm512_sub_ps( + ones, + _mm512_mul_ps(t, t)); #endif - _mm512_storeu_ps(x + i, res); + + _mm512_storeu_ps(x + j, t); + } + + i = simd_end; } #elif defined(__AVX2__) - __m256 ones = _mm256_set1_ps(1.0f); - simd_end = n - (n % 8); - for (; i < simd_end; i += 8) { - __m256 t = _mm256_loadu_ps(x + i); + const __m256 ones = _mm256_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)7; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; + j < static_cast(simd_end); + j += 8) + { + __m256 t = _mm256_loadu_ps(x + j); + t = Sleef_tanhf8_u10avx2(t); + #ifdef __FMA__ - __m256 res = _mm256_fnmadd_ps(t, t, ones); // 1 - t*t + t = _mm256_fnmadd_ps(t, t, ones); #else - __m256 res = _mm256_sub_ps(ones, _mm256_mul_ps(t, t)); + t = _mm256_sub_ps( + ones, + _mm256_mul_ps(t, t)); #endif - _mm256_storeu_ps(x + i, res); + + _mm256_storeu_ps(x + j, t); + } + + i = simd_end; } -#elif defined(__SSE4_1__) || defined(_M_AMD64) || defined(_M_X64) - __m128 ones = _mm_set1_ps(1.0f); - simd_end = n - (n % 4); - for (; i < simd_end; i += 4) +#elif defined(__AVX__) + { - __m128 t = _mm_loadu_ps(x + i); + const __m256 ones = _mm256_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)7; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; + j < static_cast(simd_end); + j += 8) + { + __m256 t = _mm256_loadu_ps(x + j); + t = Sleef_tanhf8_u10avx(t); #ifdef __FMA__ - __m128 res = _mm_fnmadd_ps(t, t, ones); // 1 - t*t + t = _mm256_fnmadd_ps(t, t, ones); #else - __m128 res = _mm_sub_ps(ones, _mm_mul_ps(t, t)); + t = _mm256_sub_ps( + ones, + _mm256_mul_ps(t, t)); #endif - _mm_storeu_ps(x + i, res); + + _mm256_storeu_ps(x + j, t); + } + + i = simd_end; } + +#elif defined(__SSE2__) || defined(_M_AMD64) || defined(_M_X64) + + { + const __m128 ones = _mm_set1_ps(1.0f); + const size_t simd_end = n & ~(size_t)3; + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; + j < static_cast(simd_end); + j += 4) + { + __m128 t = _mm_loadu_ps(x + j); + +#if defined(__SSE4_1__) + t = Sleef_tanhf4_u10sse4(t); +#else + t = Sleef_tanhf4_u10sse2(t); +#endif + +#ifdef __FMA__ + t = _mm_fnmadd_ps(t, t, ones); +#else + t = _mm_sub_ps( + ones, + _mm_mul_ps(t, t)); #endif + _mm_storeu_ps(x + j, t); + } + + i = simd_end; + } + +#else + + i = 0; + +#endif + + /* + * Scalar tail. + */ for (; i < n; ++i) - x[i] = 1.0f - x[i] * x[i]; + { + const float t = Sleef_tanhf_u10(x[i]); + x[i] = 1.0f - t * t; + } } /// @brief Two-buffer, single-pass tanh derivative: reads src, writes @@ -978,13 +1475,192 @@ namespace Deep dst[i] = sig * (1.0f - sig); } } + + static inline void d_eSigmoid(float *RESTRICT x, + const size_t n, + const bool activated = false) noexcept + { + assert(n != 0 && "n must not be 0."); + assert(x != nullptr && "x must not be null."); + + if (!activated) + e_sigmoid(x, n); + + size_t i = 0; + +#if defined(__AVX512F__) + + const __m512 two = _mm512_set1_ps(2.0f); + + const size_t r = n % 16; + const size_t simd_end = n - r; + + for (; i < simd_end; i += 16) + { + const __m512 x_512 = _mm512_loadu_ps(x + i); + + // d = 2 * x * (1 - x) + const __m512 d = _mm512_mul_ps( + _mm512_fnmadd_ps(x_512, x_512, x_512), + two); + + _mm512_storeu_ps(x + i, d); + } + +#elif defined(__AVX2__) + + const __m256 two = _mm256_set1_ps(2.0f); + + const size_t r = n % 8; + const size_t simd_end = n - r; + + for (; i < simd_end; i += 8) + { + const __m256 x_256 = _mm256_loadu_ps(x + i); + +#ifdef __FMA__ + // x * (1 - x) = x - x^2 + const __m256 d = _mm256_mul_ps( + _mm256_fnmadd_ps(x_256, x_256, x_256), + two); +#else + const __m256 d = _mm256_mul_ps( + _mm256_sub_ps(x_256, _mm256_mul_ps(x_256, x_256)), + two); +#endif + + _mm256_storeu_ps(x + i, d); + } + +#elif defined(__SSE2__) || defined(_M_AMD64) || defined(_M_X64) + + const __m128 two = _mm_set1_ps(2.0f); + + const size_t r = n % 4; + const size_t simd_end = n - r; + + for (; i < simd_end; i += 4) + { + const __m128 x_128 = _mm_loadu_ps(x + i); + +#ifdef __FMA__ + const __m128 d = _mm_mul_ps( + _mm_fnmadd_ps(x_128, x_128, x_128), + two); +#else + const __m128 d = _mm_mul_ps( + _mm_sub_ps(x_128, _mm_mul_ps(x_128, x_128)), + two); +#endif + + _mm_storeu_ps(x + i, d); + } + +#endif + + for (; i < n; ++i) + { + x[i] = 2.0f * x[i] * (1.0f - x[i]); + } + } + + static inline void d_eSigmoidInto(float *RESTRICT dst, + const float *RESTRICT src, + const size_t n) noexcept + { + assert(n != 0 && "n must not be 0."); + assert(dst != nullptr && "dst must not be null."); + assert(src != nullptr && "src must not be null."); + + size_t i = 0; + +#if defined(__AVX512F__) + + const __m512 half = _mm512_set1_ps(0.5f); + const __m512 one = _mm512_set1_ps(1.0f); + + const size_t simd_end = n - (n % 16); + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; j < (ptrdiff_t)simd_end; j += 16) + { + const __m512 s = _mm512_loadu_ps(src + j); + const __m512 a = _mm512_add_ps(_mm512_abs_ps(s), one); + const __m512 d = _mm512_div_ps(half, _mm512_mul_ps(a, a)); + + _mm512_storeu_ps(dst + j, d); + } + + i = simd_end; + +#elif defined(__AVX2__) + + const __m256 half = _mm256_set1_ps(0.5f); + const __m256 one = _mm256_set1_ps(1.0f); + const __m256 abs_mask = + _mm256_castsi256_ps(_mm256_set1_epi32(0x7FFFFFFF)); + + const size_t simd_end = n - (n % 8); + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; j < (ptrdiff_t)simd_end; j += 8) + { + const __m256 s = _mm256_loadu_ps(src + j); + const __m256 a = _mm256_add_ps( + _mm256_and_ps(s, abs_mask), + one); + + const __m256 d = _mm256_div_ps( + half, + _mm256_mul_ps(a, a)); + + _mm256_storeu_ps(dst + j, d); + } + + i = simd_end; + +#elif defined(__SSE2__) || defined(_M_AMD64) || defined(_M_X64) + + const __m128 half = _mm_set1_ps(0.5f); + const __m128 one = _mm_set1_ps(1.0f); + const __m128 abs_mask = + _mm_castsi128_ps(_mm_set1_epi32(0x7FFFFFFF)); + + const size_t simd_end = n - (n % 4); + +#pragma omp parallel for schedule(static) if (n > 16384 && !omp_in_parallel()) + for (ptrdiff_t j = 0; j < (ptrdiff_t)simd_end; j += 4) + { + const __m128 s = _mm_loadu_ps(src + j); + const __m128 a = _mm_add_ps( + _mm_and_ps(s, abs_mask), + one); + + const __m128 d = _mm_div_ps( + half, + _mm_mul_ps(a, a)); + + _mm_storeu_ps(dst + j, d); + } + + i = simd_end; + +#endif + + for (; i < n; ++i) + { + const float a = 1.0f + std::fabs(src[i]); + dst[i] = 0.5f / (a * a); + } + } + #pragma endregion static inline void linear([[maybe_unused]] float *x, [[maybe_unused]] size_t n) noexcept {} static inline void dLinear(float *x, size_t n, [[maybe_unused]] bool activated = false) noexcept { - std::fill_n(x, n, 1.0f); + std::fill(x, x + n, 1.0f); } /// @brief Two-buffer variant of dLinear -- src is unused (the @@ -992,6 +1668,6 @@ namespace Deep /// signature consistency with To_dFn2's dispatch table. static inline void dLinearInto(float *RESTRICT dst, [[maybe_unused]] const float *RESTRICT src, size_t n) noexcept { - std::fill_n(dst, n, 1.0f); + std::fill(dst, dst + n, 1.0f); } } diff --git a/include/deepity/utils/DeviceMemoryArena.h b/include/deepity/utils/DeviceMemoryArena.h index b42eaaa..00ee67c 100644 --- a/include/deepity/utils/DeviceMemoryArena.h +++ b/include/deepity/utils/DeviceMemoryArena.h @@ -1,78 +1,140 @@ #pragma once -#if defined(DEEPITY_ENABLE_CUDA) -#include +#if defined(DEEPITY_USE_CUDA) +#include +#include +#include +#include /** * @file DeviceMemoryArena.h - * @brief A simple bump-pointer allocator over a single cudaMalloc'd - * device buffer, mirroring this codebase's host-side MemoryArena but for - * GPU memory. + * @brief A 64-byte-aligned bump-pointer allocator over a single device + * buffer, mirroring this codebase's host-side MemoryArena, but backed by + * IComputeBackend::Allocate()/Free() rather than raw cudaMalloc/cudaFree + * directly -- the same allocation path Tensor already uses and + * CUDAFunctionsVerify already exercises, rather than a second, separate, + * unverified call site. * - * @warning No bounds checking: AllocateFloats() never verifies the - * running offset stays within capacity_bytes, so over-allocating past the - * arena's total_floats silently hands out a pointer past the end of the - * cudaMalloc'd buffer. Callers are responsible for summing every - * required allocation up front (mirroring how MemoryArena-based layers - * report GetRequiredFloats() before BindMemory()). + * Deliberately NOT routed through IComputeBackend on the CPU side + * (MemoryArena stays exactly as it is): CPUBackend::Allocate() has no + * huge-pages parameter, so wiring MemoryArena through it would mean + * either silently losing that real, measured feature, or bolting a + * CPU-only concept onto IComputeBackend's interface for no benefit, + * since nothing ever needs to swap MemoryArena's allocator + * polymorphically at runtime. The GPU side has no such feature to lose, + * and the virtual call only happens once per arena (construction and + * destruction), not per AllocateFloats() call -- a clean win with no + * real cost. * - * @warning cudaMalloc()'s return value is not checked in the constructor. - * If the allocation fails, base_ptr is left null/uninitialized and every - * subsequent AllocateFloats() call will silently return an invalid - * pointer rather than failing loudly. - * - * @version 1.0 - * @date 2026-06-30 - * @author Jack Rose + * @warning Individual chunks handed out by AllocateFloats() cannot be + * freed independently -- the entire arena is released at once in the + * destructor, matching MemoryArena's bump-pointer/no-reclaim design. + * @version 1.1 + * @date 2026-09-05 */ -/// @brief A bump-pointer allocator over a single contiguous block of CUDA -/// device memory. Only compiled when DEEPITY_ENABLE_CUDA is defined. -class DeviceMemoryArena +namespace Deep { -private: - /// @brief Base address of the underlying cudaMalloc'd device buffer. - float *base_ptr; - /// @brief Total capacity of the arena, in bytes. - size_t capacity_bytes; - /// @brief Current allocation offset from base_ptr, in bytes. Advances - /// monotonically with each AllocateFloats() call and is never reset - /// or reclaimed. - size_t offset_bytes; - -public: - /// @brief Allocates a single device buffer large enough to hold - /// `total_floats` floats. - /// @param total_floats Total number of floats this arena can hand - /// out across all future AllocateFloats() calls combined. - /// @warning Does not check cudaMalloc()'s return value -- see - /// file-level warning. - DeviceMemoryArena(size_t total_floats) + /// @brief A 64-byte-aligned bump-pointer allocator over a single + /// contiguous block of device memory, obtained via a supplied + /// IComputeBackend (expected to be a CUDABackend in practice, though + /// this class itself doesn't hard-require that specific type). + class DeviceMemoryArena { - capacity_bytes = total_floats * sizeof(float); - cudaMalloc(&base_ptr, capacity_bytes); - offset_bytes = 0; - } + private: + /// @brief Backend used for the underlying allocation and its + /// eventual release. Non-owning -- must outlive this arena. + IComputeBackend *backend; + /// @brief Base address of the underlying backend-allocated buffer. + float *base_ptr; + /// @brief Total capacity of the arena, in bytes, rounded up to a + /// multiple of 64. + size_t capacity_bytes; + /// @brief Current allocation offset from base_ptr, in bytes. + /// Advances monotonically with each AllocateFloats() call and is + /// never reset or reclaimed. + size_t offset_bytes; - /// @brief Frees the underlying device buffer via cudaFree(). - ~DeviceMemoryArena() - { - cudaFree(base_ptr); - } + public: + /// @brief Allocates a single 64-byte-aligned device buffer large + /// enough to hold `total_floats` floats, rounded up to the + /// nearest 64-byte boundary. + /// @param backend The backend to allocate from (and later free + /// through). Must outlive this DeviceMemoryArena. Passing a + /// CPUBackend here would allocate host memory under a class + /// named "device" -- that's a caller contract, not something + /// this class checks at runtime, mirroring Tensor's own + /// backend/device pairing contract. + /// @param total_floats Total number of floats this arena can + /// hand out across all future AllocateFloats() calls combined. + /// @throws std::bad_alloc if the underlying allocation fails. + DeviceMemoryArena(IComputeBackend *backend, size_t total_floats) + : backend(backend) + { + capacity_bytes = (total_floats * sizeof(float) + 63) & ~(size_t)63; - /// @brief Hands out a chunk of `num_floats` floats from the arena via - /// simple pointer-bump allocation. - /// @param num_floats Number of floats to allocate from the arena. - /// @return Pointer to the start of the allocated chunk, valid for the - /// lifetime of this DeviceMemoryArena. - /// @warning No bounds checking against capacity_bytes -- see - /// file-level warning. Individual chunks are also never freed - /// independently; the entire arena is released at once in the - /// destructor. - float *AllocateFloats(size_t num_floats) - { - float *chunk = base_ptr + (offset_bytes / sizeof(float)); - offset_bytes += num_floats * sizeof(float); - return chunk; - } -}; + base_ptr = backend->Allocate(capacity_bytes / sizeof(float)); + if (!base_ptr) + { + throw std::bad_alloc(); + } + offset_bytes = 0; + } + + /// @brief Frees the underlying device buffer via the same + /// backend it was allocated from. + ~DeviceMemoryArena() + { + if (base_ptr) + { + backend->Free(base_ptr); + } + } + + // Delete copy/move constructors to prevent double-free corruption + // -- the original version of this class had no such guard, a real + // gap relative to MemoryArena's own established protection. + + /// @brief Deleted: DeviceMemoryArena owns a single device + /// allocation, so copying would risk a double-free. + DeviceMemoryArena(const DeviceMemoryArena &) = delete; + /// @brief Deleted: DeviceMemoryArena owns a single device + /// allocation, so copy-assignment would risk a double-free. + DeviceMemoryArena &operator=(const DeviceMemoryArena &) = delete; + + /// @brief Allocates a 64-byte aligned chunk of floats from the + /// arena. + /// @param num_floats Number of floats to allocate from the + /// arena. The actual reservation is rounded up to the nearest + /// 64-byte boundary, matching MemoryArena's own guarantee (the + /// original version of this class had no such rounding at all). + /// @return Pointer to the start of the allocated chunk, valid + /// for the lifetime of this DeviceMemoryArena. 64-byte aligned. + /// @throws std::runtime_error if the requested allocation would + /// exceed the arena's total capacity -- the original version of + /// this class had no bounds checking at all. + /// @warning Individual chunks are never freed independently; the + /// entire arena is released at once in the destructor. + float *AllocateFloats(size_t num_floats) + { + size_t allocation_size = (num_floats * sizeof(float) + 63) & ~(size_t)63; + + if (offset_bytes + allocation_size > capacity_bytes) + { + throw std::runtime_error("Fatal: DeviceMemoryArena capacity exceeded during allocation."); + } + + float *chunk = reinterpret_cast( + reinterpret_cast(base_ptr) + offset_bytes); + + offset_bytes += allocation_size; + return chunk; + } + + /// @brief Returns how many bytes have been allocated from the + /// arena so far. + size_t GetUsedBytes() const { return offset_bytes; } + /// @brief Returns the arena's total capacity. + size_t GetCapacityBytes() const { return capacity_bytes; } + }; +} #endif \ No newline at end of file diff --git a/logs/build.log b/logs/build.log index ca3e1cf..b29064d 100644 --- a/logs/build.log +++ b/logs/build.log @@ -1,23 +1,41 @@ ---- Deepity Build Log (Release, arch=native) --- +--- Deepity Build Log (Release, arch=fast) --- === CMake Configuration === --- PGO: instrumented (GENERATE) build -- profile data will be written to /home/rose0/Projects/deepity/build/pgo-data +-- The C compiler identification is GNU 16.2.1 +-- The CXX compiler identification is GNU 16.2.1 +-- Detecting C compiler ABI info +-- Detecting C compiler ABI info - done +-- Check for working C compiler: /usr/bin/cc - skipped +-- Detecting C compile features +-- Detecting C compile features - done +-- Detecting CXX compiler ABI info +-- Detecting CXX compiler ABI info - done +-- Check for working CXX compiler: /usr/bin/c++ - skipped +-- Detecting CXX compile features +-- Detecting CXX compile features - done -- CXX compiler: /usr/bin/c++ -- Found OpenMP_C: -fopenmp (found version "5.2") -- Found OpenMP_CXX: -fopenmp (found version "5.2") --- Using Intel MKL (found via MKLConfig.cmake). +-- Found OpenMP: TRUE (found version "5.2") +-- Found Python: /home/rose0/Projects/deepity/.venv/bin/python3.14 (found suitable version "3.14.7", minimum required is "3.8") found components: Interpreter Development.Module +-- Performing Test NB_HAS_MTLS_GNU2 +-- Performing Test NB_HAS_MTLS_GNU2 - Success +-- Looking for sgemm_ +-- Looking for sgemm_ - found +-- Found BLAS: /usr/lib64/libopenblas.so -- Configuring SLEEF 3.9.0 -- Could NOT find OpenSSL, try to set the path to OpenSSL root folder in the system variable OPENSSL_ROOT_DIR (missing: OPENSSL_CRYPTO_LIBRARY OPENSSL_INCLUDE_DIR) +-- Found PkgConfig: /usr/bin/pkg-config (found version "2.5.1") -- Found OpenMP_C: -fopenmp (found version "5.2") -- Found OpenMP_CXX: -fopenmp (found version "5.2") -- Configuring build for SLEEF-v3.9.0 - Target system: Linux-7.1.12-200.fc44.x86_64 + Target system: Linux-7.2.4-200.fc44.x86_64 Target processor: x86_64 - Host system: Linux-7.1.12-200.fc44.x86_64 + Host system: Linux-7.2.4-200.fc44.x86_64 Host processor: x86_64 Detected C compiler: GNU @ /usr/bin/cc CMake: 4.3.0 - Make program: /usr/bin/gmake + Make program: /usr/bin/ninja-build CMake build type: Release -- Using option `-Wall -Wno-unused-function -Wno-attributes -Wno-unused-result -Wno-psabi -ffp-contract=off -fno-math-errno -fno-trapping-math` to compile libsleef -- Building shared libs : OFF @@ -31,296 +49,346 @@ -- SDE : SDE_COMMAND-NOTFOUND -- COMPILER_SUPPORTS_OPENMP : 1 -- A version of SLEEF compatible with libm and libmvec in GNU libc will be produced (sleefgnuabi.so) --- Failed to find LLVM FileCheck --- Google Benchmark version: v1.9.1, normalized to 1.9.1 --- Performing Test HAVE_STD_REGEX -- success --- Performing Test HAVE_GNU_POSIX_REGEX -- failed to compile --- Performing Test HAVE_POSIX_REGEX -- success --- Performing Test HAVE_STEADY_CLOCK -- success --- Performing Test HAVE_PTHREAD_AFFINITY -- success --- Configuring done (3.0s) --- Generating done (0.3s) +-- Configuring done (5.0s) +-- Generating done (0.0s) -- Build files have been written to: /home/rose0/Projects/deepity/build/Release === Compilation === -[ 1%] Building C object _deps/sleef-build/src/common/CMakeFiles/common.dir/common.c.o -[ 1%] Building C object _deps/sleef-build/src/common/CMakeFiles/addSuffix.dir/addSuffix.c.o -[ 1%] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkrename_gnuabi.dir/mkrename_gnuabi.c.o -[ 1%] Building C object _deps/sleef-build/src/common/CMakeFiles/qtesterutil_obj.dir/qtesterutil.c.o -[ 1%] Building CXX object _deps/sleef-build/src/common/CMakeFiles/psha_obj.dir/psha2_capi.cpp.o -[ 1%] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkdisp.dir/mkdisp.c.o -[ 1%] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkmasked_gnuabi.dir/mkmasked_gnuabi.c.o -[ 1%] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkrename.dir/mkrename.c.o -[ 2%] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkalias.dir/mkalias.c.o -[ 3%] Building C object _deps/sleef-build/src/common/CMakeFiles/testerutil_obj.dir/testerutil.c.o -[ 4%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_internals.cpp.o -[ 5%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_type.cpp.o -[ 5%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/benchmark_api_internal.cc.o -[ 6%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/benchmark.cc.o -[ 6%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_func.cpp.o -[ 5%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_enum.cpp.o -[ 6%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_static_property.cpp.o -[ 7%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/benchmark_name.cc.o -[ 8%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_ndarray.cpp.o -[ 8%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/benchmark_register.cc.o -[ 8%] Built target common -[ 8%] Linking C executable ../../bin/mkrename_gnuabi -[ 9%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/benchmark_runner.cc.o -[ 9%] Linking C executable ../../bin/mkalias -[ 9%] Built target mkrename_gnuabi -[ 9%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/check.cc.o -[ 10%] Linking C executable ../../bin/mkdisp -[ 10%] Built target qtesterutil_obj -[ 11%] Linking C executable ../../bin/mkmasked_gnuabi -[ 11%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/colorprint.cc.o -[ 11%] Linking C executable ../../bin/addSuffix -[ 11%] Built target mkalias -[ 11%] Built target testerutil_obj -[ 11%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_datetime.cpp.o -[ 11%] Built target mkdisp -[ 12%] Generating include/renamesse2_gnuabi.h -Generating renamesse2_gnuabi.h: mkrename_gnuabi sse2 b 2 4 _mm128d _mm128 _mm128i _mm128i __SSE2__ -[ 13%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/common.cpp.o -[ 13%] Built target mkmasked_gnuabi -[ 14%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/commandlineflags.cc.o -[ 14%] Built target addSuffix -[ 14%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabisse2dp.dir/sleefsimddp.c.o -[ 14%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/complexity.cc.o -[ 15%] Linking C executable ../../bin/mkrename -[ 15%] Built target mkrename -[ 15%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabisse2sp.dir/sleefsimdsp.c.o -[ 15%] Built target psha_obj -[ 15%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/error.cpp.o -[ 16%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/trampoline.cpp.o -[ 16%] Generating include/renameavx_gnuabi.h -Generating renameavx_gnuabi.h: mkrename_gnuabi avx c 4 8 __m256d __m256 __m128i struct\ {\ __m128i\ x,\ y;\ } __AVX__ -[ 16%] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/implicit.cpp.o -[ 16%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavxdp.dir/sleefsimddp.c.o -[ 16%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavxsp.dir/sleefsimdsp.c.o -[ 17%] Generating include/renameavx2_gnuabi.h +[1/152] Building C object _deps/sleef-build/src/common/CMakeFiles/common.dir/common.c.o +[2/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkrename_gnuabi.dir/mkrename_gnuabi.c.o +[3/152] Linking C executable _deps/sleef-build/bin/mkrename_gnuabi +[4/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkalias.dir/mkalias.c.o +[5/152] Generating include/renameavx2_gnuabi.h Generating renameavx2_gnuabi.h: mkrename_gnuabi avx2 d 4 8 __m256d __m256 __m128i __m256i __AVX2__ -[ 18%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/console_reporter.cc.o -[ 18%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx2dp.dir/sleefsimddp.c.o -[ 19%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx2sp.dir/sleefsimdsp.c.o -[ 19%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/counter.cc.o -[ 19%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/csv_reporter.cc.o -[ 20%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/json_reporter.cc.o -[ 20%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/perf_counters.cc.o -[ 21%] Generating include/masked_avx512f_dp_gnuabi.h -[ 21%] Generating include/masked_avx512f_sp_gnuabi.h -[ 21%] Built target maskedAVX512F_generated -[ 21%] Generating sleeflibm_SSE_.h.tmp -[ 22%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/reporter.cc.o -[ 22%] Generating sleeflibm_SSE2.h.tmp -[ 23%] Generating sleeflibm_SSE4.h.tmp -[ 23%] Generating sleeflibm_AVX_.h.tmp -[ 24%] Generating sleeflibm_AVX.h.tmp -[ 24%] Generating include/renameavx512fnofma.h +[6/152] Generating include/renameavx512f_gnuabi.h +Generating renameavx512f_gnuabi.h: mkrename_gnuabi avx512f e 8 16 __m512d __m512 __m256i __m512i __AVX512F__ +[7/152] Generating include/renameavx_gnuabi.h +Generating renameavx_gnuabi.h: mkrename_gnuabi avx c 4 8 __m256d __m256 __m128i struct\ {\ __m128i\ x,\ y;\ } __AVX__ +[8/152] Generating include/renamesse2_gnuabi.h +Generating renamesse2_gnuabi.h: mkrename_gnuabi sse2 b 2 4 _mm128d _mm128 _mm128i _mm128i __SSE2__ +[9/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkdisp.dir/mkdisp.c.o +[10/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkmasked_gnuabi.dir/mkmasked_gnuabi.c.o +[11/152] Building C object _deps/sleef-build/src/common/CMakeFiles/addSuffix.dir/addSuffix.c.o +[12/152] Linking C executable _deps/sleef-build/bin/mkdisp +[13/152] Generating dispscalar.c.body +[14/152] Linking C executable _deps/sleef-build/bin/mkalias +[15/152] Generating dispsse.c.tmp +[16/152] Generating alias_AVX512F_sp.h.tmp +[17/152] Generating alias_AVX512F_dp.h.tmp +[18/152] Building C object _deps/sleef-build/src/common/CMakeFiles/qtesterutil_obj.dir/qtesterutil.c.o +[19/152] Generating dispavx.c.tmp +[20/152] Building C object _deps/sleef-build/src/common/CMakeFiles/testerutil_obj.dir/testerutil.c.o +[21/152] Generating include/alias_avx512f.h +[22/152] Linking C executable _deps/sleef-build/bin/mkmasked_gnuabi +[23/152] Generating dispscalar.c +[24/152] Generating dispsse.c +[25/152] Generating include/masked_avx512f_dp_gnuabi.h +[26/152] Generating include/masked_avx512f_sp_gnuabi.h +[27/152] Generating dispavx.c +[28/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/mkrename.dir/mkrename.c.o +[29/152] Linking C executable _deps/sleef-build/bin/mkrename +[30/152] Generating sleeflibm_AVX.h.tmp +[31/152] Generating sleeflibm_AVX2.h.tmp +[32/152] Generating sleeflibm_AVX2128.h.tmp +[33/152] Generating sleeflibm_AVX512F.h.tmp +[34/152] Generating sleeflibm_AVX512FNOFMA.h.tmp +[35/152] Generating sleeflibm_AVX512F_.h.tmp +[36/152] Generating sleeflibm_AVX_.h.tmp +[37/152] Generating sleeflibm_DSP_SCALAR.h.tmp +[38/152] Generating sleeflibm_FMA4.h.tmp +[39/152] Generating sleeflibm_PURECFMA_SCALAR.h.tmp +[40/152] Generating sleeflibm_PUREC_SCALAR.h.tmp +[41/152] Generating sleeflibm_SSE2.h.tmp +[42/152] Generating sleeflibm_SSE4.h.tmp +[43/152] Generating sleeflibm_SSE_.h.tmp +[44/152] Generating include/renameavx512fnofma.h Generating renameavx512fnofma.h: mkrename cinz_ 8 16 avx512fnofma -[ 24%] Generating sleeflibm_FMA4.h.tmp -[ 24%] Built target renameAVX512FNOFMA.h_generated -[ 24%] Generating sleeflibm_AVX2.h.tmp -[ 25%] Generating sleeflibm_AVX2128.h.tmp -[ 25%] Generating alias_AVX512F_dp.h.tmp -[ 26%] Generating sleeflibm_AVX512F_.h.tmp -[ 26%] Generating alias_AVX512F_sp.h.tmp -[ 26%] Generating sleeflibm_AVX512F.h.tmp -[ 27%] Generating include/alias_avx512f.h -[ 27%] Generating sleeflibm_AVX512FNOFMA.h.tmp -[ 28%] Built target sleefgnuabisse2sp -[ 28%] Built target alias_avx512f.h_generated -[ 28%] Generating sleeflibm_PUREC_SCALAR.h.tmp -[ 28%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/statistics.cc.o -[ 29%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/string_util.cc.o -[ 30%] Generating sleeflibm_PURECFMA_SCALAR.h.tmp -[ 31%] Generating sleeflibm_DSP_SCALAR.h.tmp -[ 31%] Generating include/renameavx512f.h +[45/152] Generating include/renameavx512f.h Generating renameavx512f.h: mkrename finz_ 8 16 avx512f -[ 31%] Built target sleefgnuabisse2dp -[ 32%] Generating include/renameavx2.h +[46/152] Generating include/renameavx2.h Generating renameavx2.h: mkrename finz_ 4 8 avx2 -[ 32%] Generating include/renameavx2128.h +[47/152] Generating include/renameavx2128.h Generating renameavx2128.h: mkrename finz_ 2 4 avx2128 -[ 32%] Built target renameAVX512F.h_generated -[ 32%] Generating include/renamefma4.h +[48/152] Generating include/renamefma4.h Generating renamefma4.h: mkrename finz_ 4 8 fma4 -[ 32%] Generating include/renameavx.h -[ 33%] Built target renameAVX2128.h_generated +[49/152] Generating include/renameavx.h Generating renameavx.h: mkrename cinz_ 4 8 avx -[ 33%] Built target renameAVX2.h_generated -[ 33%] Generating include/renamesse4.h +[50/152] Generating include/renamesse4.h Generating renamesse4.h: mkrename cinz_ 2 4 sse4 -[ 33%] Built target renameFMA4.h_generated -[ 34%] Generating include/renamesse2.h +[51/152] Generating include/renamesse2.h Generating renamesse2.h: mkrename cinz_ 2 4 sse2 -[ 35%] Built target renameAVX.h_generated -[ 35%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/sysinfo.cc.o -[ 35%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark.dir/timers.cc.o -[ 36%] Generating include/renamepurec_scalar.h +[52/152] Generating include/renamepurec_scalar.h Generating renamepurec_scalar.h: mkrename cinz_ 1 1 purec -[ 36%] Generating include/renamepurecfma_scalar.h +[53/152] Generating include/renamepurecfma_scalar.h Generating renamepurecfma_scalar.h: mkrename finz_ 1 1 purecfma -[ 37%] Generating include/renamecuda.h +[54/152] Generating include/renamecuda.h Generating renamecuda.h: mkrename finz_ 1 1 cuda -[ 38%] Built target renameSSE4.h_generated -[ 39%] Generating include/renameavx512f_gnuabi.h -Generating renameavx512f_gnuabi.h: mkrename_gnuabi avx512f e 8 16 __m512d __m512 __m256i __m512i __AVX512F__ -[ 39%] Generating ../../include/sleef.h -[ 39%] Built target renameSSE2.h_generated -[ 40%] Built target headers -[ 40%] Built target renamePUREC_SCALAR.h_generated -[ 41%] Built target renamePURECFMA_SCALAR.h_generated -[ 41%] Generating include/renamedspscalar.h -[ 41%] Generating dispscalar.c.body -[ 41%] Built target renamedspscalar.h_generated -[ 42%] Generating dispscalar.c -[ 42%] Generating include/renamedsp128.h -[ 42%] Built target dispscalar.c_generated -[ 42%] Built target renamedsp128.h_generated -[ 42%] Generating dispsse.c.tmp -[ 43%] Generating include/renamedsp256.h -[ 43%] Built target sleefgnuabiavx2dp -[ 43%] Built target renamedsp256.h_generated -[ 44%] Generating dispsse.c -[ 44%] Generating dispavx.c.tmp -[ 44%] Built target dispsse.c_generated -[ 45%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx512fdp.dir/sleefsimddp.c.o -[ 46%] Generating dispavx.c -[ 47%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx512fsp.dir/sleefsimdsp.c.o -[ 47%] Built target dispavx.c_generated -[ 48%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512fnofma.dir/sleefsimdsp.c.o -[ 48%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512fnofma.dir/sleefsimdsp.c.o -[ 48%] Built target sleefgnuabiavxdp -[ 49%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512fnofma.dir/sleefsimddp.c.o -[ 49%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512fnofma.dir/sleefsimddp.c.o -[ 50%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512f.dir/sleefsimdsp.c.o -[ 50%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512f.dir/sleefsimddp.c.o -[ 50%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512f.dir/sleefsimdsp.c.o -[ 50%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2.dir/sleefsimdsp.c.o -[ 50%] Built target sleefgnuabiavx2sp -[ 50%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512f.dir/sleefsimddp.c.o -[ 51%] Built target sleefgnuabiavxsp -[ 51%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2.dir/sleefsimddp.c.o -[ 51%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2.dir/sleefsimdsp.c.o -[ 51%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2128.dir/sleefsimdsp.c.o -[ 52%] Linking CXX static library bin/libnanobind-static.a -[ 52%] Built target nanobind-static -[ 53%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2.dir/sleefsimddp.c.o -[ 54%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2128.dir/sleefsimddp.c.o -[ 55%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2128.dir/sleefsimdsp.c.o -[ 56%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleeffma4.dir/sleefsimdsp.c.o -[ 57%] Built target sleefdetavx512fnofma -[ 58%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetfma4.dir/sleefsimdsp.c.o -[ 58%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleeffma4.dir/sleefsimddp.c.o -[ 59%] Built target sleefdetavx512f -[ 59%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetfma4.dir/sleefsimddp.c.o -[ 59%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2128.dir/sleefsimddp.c.o -[ 60%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx.dir/sleefsimdsp.c.o -[ 60%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx.dir/sleefsimdsp.c.o -[ 60%] Built target sleefdetavx2 -[ 61%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx.dir/sleefsimddp.c.o -[ 61%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx.dir/sleefsimddp.c.o -[ 61%] Built target sleefavx512fnofma -[ 61%] Built target sleefgnuabiavx512fdp -[ 61%] Built target sleefgnuabiavx512fsp -[ 61%] Built target sleefavx512f -[ 61%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse2.dir/sleefsimdsp.c.o -[ 61%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse4.dir/sleefsimddp.c.o -[ 62%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse4.dir/sleefsimdsp.c.o -[ 63%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse2.dir/sleefsimddp.c.o -[ 64%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse4.dir/sleefsimdsp.c.o -[ 64%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse2.dir/sleefsimdsp.c.o -[ 65%] Built target sleefavx2 -[ 65%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse4.dir/sleefsimddp.c.o -[ 66%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse2.dir/sleefsimddp.c.o -[ 66%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurec_scalar.dir/sleefsimddp.c.o -[ 67%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurec_scalar.dir/sleefsimdsp.c.o -[ 67%] Built target sleefdetavx2128 -[ 68%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurec_scalar.dir/sleefsimdsp.c.o -[ 68%] Built target sleefdetfma4 -[ 68%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurecfma_scalar.dir/sleefsimdsp.c.o -[ 69%] Built target sleefavx2128 -[ 70%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurecfma_scalar.dir/sleefsimddp.c.o -[ 70%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurec_scalar.dir/sleefsimddp.c.o -[ 70%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurecfma_scalar.dir/sleefsimdsp.c.o -[ 70%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurecfma_scalar.dir/sleefsimddp.c.o -[ 70%] Building C object _deps/sleef-build/src/libm/CMakeFiles/dispscalar_obj.dir/dispscalar.c.o -[ 71%] Building C object _deps/sleef-build/src/libm/CMakeFiles/dispsse_obj.dir/dispsse.c.o -[ 72%] Built target sleefdetavx -[ 72%] Building C object _deps/sleef-build/src/libm/CMakeFiles/dispavx_obj.dir/dispavx.c.o -[ 72%] Built target sleefdetsse4 -[ 73%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabi.dir/rempitab.c.o -[ 73%] Linking C static library ../../lib/libsleefgnuabi.a -[ 73%] Built target sleeffma4 -[ 74%] Built target sleefdetsse2 -[ 74%] Built target sleefgnuabi -[ 75%] Built target sleefdetpurecfma_scalar -[ 75%] Built target sleefdetpurec_scalar -[ 75%] Built target sleefsse2 -[ 75%] Built target sleefsse4 -[ 75%] Built target sleefavx -[ 76%] Built target sleefpurecfma_scalar -[ 76%] Built target sleefpurec_scalar -[ 77%] Built target dispscalar_obj -[ 78%] Built target dispsse_obj -[ 79%] Built target dispavx_obj -[ 79%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleef.dir/sleefld.c.o -[ 79%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleef.dir/sleefqp.c.o -[ 80%] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleef.dir/rempitab.c.o -[ 81%] Linking C static library ../../lib/libsleef.a -[ 81%] Built target sleef -[ 81%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/SimplePCLayer.cpp.o -[ 81%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/DiscriminativePCLayer.cpp.o -[ 82%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/DiscriminativePCNetwork.cpp.o -[ 83%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/RBLayer.cpp.o -[ 83%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/ConvPCLayer.cpp.o -[ 83%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/ConvPCNetwork.cpp.o -[ 84%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/SimpleConvPCLayer.cpp.o -[ 84%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/StreamAlignedBatcher.cpp.o -[ 84%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/SimplePCNetwork.cpp.o -[ 85%] Building CXX object CMakeFiles/DeepityProfiled.dir/src/ModelIO.cpp.o -[ 86%] Building CXX object CMakeFiles/Deepity.dir/src/DiscriminativePCLayer.cpp.o -[ 86%] Building CXX object CMakeFiles/Deepity.dir/src/RBLayer.cpp.o -[ 86%] Building CXX object CMakeFiles/Deepity.dir/src/SimpleConvPCLayer.cpp.o -[ 87%] Building CXX object CMakeFiles/Deepity.dir/src/ConvPCLayer.cpp.o -[ 88%] Building CXX object CMakeFiles/Deepity.dir/src/GaussSeidelPCLayer.cpp.o -[ 88%] Building CXX object CMakeFiles/Deepity.dir/src/SimplePCLayer.cpp.o -[ 88%] Building CXX object CMakeFiles/Deepity.dir/src/DirectKPPCLayer.cpp.o -[ 88%] Building CXX object CMakeFiles/Deepity.dir/src/ConvPCNetwork.cpp.o -[ 89%] Building CXX object CMakeFiles/Deepity.dir/src/DiscriminativePCNetwork.cpp.o -[ 90%] Linking CXX static library ../../../bin/libbenchmark.a -[ 90%] Built target benchmark -[ 91%] Building CXX object CMakeFiles/Deepity.dir/src/SimplePCNetwork.cpp.o -[ 91%] Building CXX object CMakeFiles/Deepity.dir/src/GaussSeidelPCNetwork.cpp.o -[ 91%] Building CXX object _deps/benchmark-build/src/CMakeFiles/benchmark_main.dir/benchmark_main.cc.o -[ 91%] Building CXX object CMakeFiles/Deepity.dir/src/SimpleConvPCNetwork.cpp.o -[ 92%] Building CXX object CMakeFiles/Deepity.dir/src/DirectKPPCNetwork.cpp.o -[ 92%] Building CXX object CMakeFiles/Deepity.dir/src/ModelIO.cpp.o -[ 93%] Building CXX object CMakeFiles/Deepity.dir/src/StreamAlignedBatcher.cpp.o -[ 94%] Linking CXX static library ../../../bin/libbenchmark_main.a -[ 94%] Built target benchmark_main -[ 95%] Linking CXX static library bin/libDeepityProfiled.a -[ 95%] Built target DeepityProfiled -[ 95%] Building CXX object CMakeFiles/ActivationBenchmark.dir/tests/tActivations.cpp.o -[ 97%] Building CXX object CMakeFiles/DeepityProfile.dir/tests/tProfile.cpp.o -[ 97%] Building CXX object CMakeFiles/GflopsBenchmark.dir/tests/tGflopsBenchmark.cpp.o -[ 97%] Linking CXX static library bin/libDeepity.a -[ 97%] Built target Deepity -[ 97%] Building CXX object CMakeFiles/GaussSeidelMiddleLayerVerify.dir/tests/tGaussSeidelMiddleLayerVerify.cpp.o -[ 97%] Building CXX object CMakeFiles/pydeepity.dir/bindings/pybinding.cpp.o -[ 97%] Building CXX object CMakeFiles/DirectKPVerify.dir/tests/tDirectKPVerify.cpp.o -[ 98%] Linking CXX executable bin/DirectKPVerify -[ 98%] Linking CXX executable bin/DeepityProfile -[ 98%] Linking CXX executable bin/GflopsBenchmark -[ 98%] Linking CXX executable bin/GaussSeidelMiddleLayerVerify -[ 98%] Built target DeepityProfile -[ 98%] Built target GflopsBenchmark -[ 98%] Linking CXX executable bin/ActivationBenchmark -[ 98%] Built target ActivationBenchmark -[ 98%] Built target GaussSeidelMiddleLayerVerify -[ 98%] Built target DirectKPVerify -[100%] Linking CXX shared module bin/pydeepity.cpython-314-x86_64-linux-gnu.so -[100%] Built target pydeepity +[55/152] Generating ../../include/sleef.h +[56/152] Generating include/renamedspscalar.h +[57/152] Generating include/renamedsp128.h +[58/152] Generating include/renamedsp256.h +[59/152] Building CXX object _deps/sleef-build/src/common/CMakeFiles/psha_obj.dir/psha2_capi.cpp.o +[60/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_static_property.cpp.o +[61/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/implicit.cpp.o +[62/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/trampoline.cpp.o +[63/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_datetime.cpp.o +[64/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/error.cpp.o +[65/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx2sp.dir/sleefsimdsp.c.o +[66/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/common.cpp.o +[67/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_ndarray.cpp.o +[68/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx2dp.dir/sleefsimddp.c.o +[69/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512fnofma.dir/sleefsimdsp.c.o +[70/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512fnofma.dir/sleefsimddp.c.o +[71/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512f.dir/sleefsimdsp.c.o +[72/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabisse2sp.dir/sleefsimdsp.c.o +[73/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx512fdp.dir/sleefsimddp.c.o +[74/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabisse2dp.dir/sleefsimddp.c.o +[75/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512fnofma.dir/sleefsimdsp.c.o +[76/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512fnofma.dir/sleefsimddp.c.o +[77/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavxdp.dir/sleefsimddp.c.o +[78/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_internals.cpp.o +[79/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2.dir/sleefsimdsp.c.o +[80/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavxsp.dir/sleefsimdsp.c.o +[81/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512f.dir/sleefsimdsp.c.o +[82/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_enum.cpp.o +[83/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx512f.dir/sleefsimddp.c.o +[84/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2128.dir/sleefsimdsp.c.o +[85/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx512f.dir/sleefsimddp.c.o +[86/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2.dir/sleefsimddp.c.o +[87/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_func.cpp.o +[88/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx2128.dir/sleefsimddp.c.o +[89/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2.dir/sleefsimdsp.c.o +[90/152] Building CXX object CMakeFiles/nanobind-static.dir/.venv/lib/python3.14/site-packages/nanobind/src/nb_type.cpp.o +[91/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetfma4.dir/sleefsimdsp.c.o +[92/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2.dir/sleefsimddp.c.o +[93/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2128.dir/sleefsimdsp.c.o +[94/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetfma4.dir/sleefsimddp.c.o +[95/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx2128.dir/sleefsimddp.c.o +[96/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx.dir/sleefsimdsp.c.o +[97/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse4.dir/sleefsimdsp.c.o +[98/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetavx.dir/sleefsimddp.c.o +[99/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurec_scalar.dir/sleefsimdsp.c.o +[100/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleef.dir/rempitab.c.o +[101/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleef.dir/sleefld.c.o +[102/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse4.dir/sleefsimddp.c.o +[103/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleeffma4.dir/sleefsimdsp.c.o +[104/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse2.dir/sleefsimdsp.c.o +[105/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurecfma_scalar.dir/sleefsimdsp.c.o +[106/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleef.dir/sleefqp.c.o +[107/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleeffma4.dir/sleefsimddp.c.o +[108/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetsse2.dir/sleefsimddp.c.o +[109/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurec_scalar.dir/sleefsimddp.c.o +[110/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefdetpurecfma_scalar.dir/sleefsimddp.c.o +[111/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse4.dir/sleefsimdsp.c.o +[112/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx.dir/sleefsimddp.c.o +[113/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse4.dir/sleefsimddp.c.o +[114/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefavx.dir/sleefsimdsp.c.o +[115/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse2.dir/sleefsimdsp.c.o +[116/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurecfma_scalar.dir/sleefsimdsp.c.o +[117/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/dispscalar_obj.dir/dispscalar.c.o +[118/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurecfma_scalar.dir/sleefsimddp.c.o +[119/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurec_scalar.dir/sleefsimdsp.c.o +[120/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefsse2.dir/sleefsimddp.c.o +[121/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefpurec_scalar.dir/sleefsimddp.c.o +[122/152] Building CXX object CMakeFiles/Deepity.dir/src/RBLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/RBLayer.h:4, + from /home/rose0/Projects/deepity/src/RBLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[123/152] Building CXX object CMakeFiles/Deepity.dir/src/ConvPCLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/ConvPCLayer.h:8, + from /home/rose0/Projects/deepity/src/ConvPCLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[124/152] Building CXX object CMakeFiles/Deepity.dir/src/StreamAlignedBatcher.cpp.o +[125/152] Linking CXX static library bin/libnanobind-static.a +[126/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/dispavx_obj.dir/dispavx.c.o +[127/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabi.dir/rempitab.c.o +[128/152] Linking C executable _deps/sleef-build/bin/addSuffix +[129/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/dispsse_obj.dir/dispsse.c.o +[130/152] Building CXX object CMakeFiles/Deepity.dir/src/GaussSeidelPCLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/GaussSeidelPCLayer.h:8, + from /home/rose0/Projects/deepity/src/GaussSeidelPCLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[131/152] Building CXX object CMakeFiles/Deepity.dir/src/SimpleConvPCLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/SimpleConvPCLayer.h:8, + from /home/rose0/Projects/deepity/src/SimpleConvPCLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[132/152] Building CXX object CMakeFiles/Deepity.dir/src/DirectKPPCLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/DirectKPPCLayer.h:6, + from /home/rose0/Projects/deepity/src/DirectKPPCLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[133/152] Linking C static library _deps/sleef-build/lib/libsleef.a +[134/152] Building CXX object CMakeFiles/Deepity.dir/src/DiscriminativePCLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/DiscriminativePCLayer.h:8, + from /home/rose0/Projects/deepity/src/DiscriminativePCLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[135/152] Building CXX object CMakeFiles/Deepity.dir/src/DiscriminativePCNetwork.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/DiscriminativePCLayer.h:8, + from /home/rose0/Projects/deepity/include/deepity/networks/DiscriminativePCNetwork.h:5, + from /home/rose0/Projects/deepity/src/DiscriminativePCNetwork.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[136/152] Building CXX object CMakeFiles/Deepity.dir/src/backend/Tensor.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/backend/IComputeBackend.h:3, + from /home/rose0/Projects/deepity/include/deepity/backend/Tensor.h:4, + from /home/rose0/Projects/deepity/src/backend/Tensor.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[137/152] Building CXX object CMakeFiles/Deepity.dir/src/SimplePCNetwork.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/SimplePCLayer.h:8, + from /home/rose0/Projects/deepity/include/deepity/networks/SimplePCNetwork.h:5, + from /home/rose0/Projects/deepity/src/SimplePCNetwork.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[138/152] Building CXX object CMakeFiles/Deepity.dir/src/backend/Backend.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/backend/IComputeBackend.h:3, + from /home/rose0/Projects/deepity/include/deepity/backend/Backend.h:3, + from /home/rose0/Projects/deepity/src/backend/Backend.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[139/152] Building CXX object CMakeFiles/Deepity.dir/src/ConvPCNetwork.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/ConvPCLayer.h:8, + from /home/rose0/Projects/deepity/include/deepity/networks/ConvPCNetwork.h:5, + from /home/rose0/Projects/deepity/src/ConvPCNetwork.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[140/152] Building CXX object CMakeFiles/Deepity.dir/src/GaussSeidelPCNetwork.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/GaussSeidelPCLayer.h:8, + from /home/rose0/Projects/deepity/include/deepity/networks/GaussSeidelPCNetwork.h:5, + from /home/rose0/Projects/deepity/src/GaussSeidelPCNetwork.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[141/152] Building CXX object CMakeFiles/Deepity.dir/src/SimplePCLayer.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/SimplePCLayer.h:8, + from /home/rose0/Projects/deepity/src/SimplePCLayer.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[142/152] Building CXX object CMakeFiles/Deepity.dir/src/backend/CPUBackend.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/backend/IComputeBackend.h:3, + from /home/rose0/Projects/deepity/include/deepity/backend/CPUBackend.h:2, + from /home/rose0/Projects/deepity/src/backend/CPUBackend.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[143/152] Building CXX object CMakeFiles/Deepity.dir/src/SimpleConvPCNetwork.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/SimpleConvPCLayer.h:8, + from /home/rose0/Projects/deepity/include/deepity/networks/SimpleConvPCNetwork.h:6, + from /home/rose0/Projects/deepity/src/SimpleConvPCNetwork.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[144/152] Building CXX object CMakeFiles/Deepity.dir/src/DirectKPPCNetwork.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/DirectKPPCLayer.h:6, + from /home/rose0/Projects/deepity/include/deepity/networks/DirectKPPCNetwork.h:5, + from /home/rose0/Projects/deepity/src/DirectKPPCNetwork.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[145/152] Building CXX object CMakeFiles/Deepity.dir/src/ModelIO.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/layers/DiscriminativePCLayer.h:8, + from /home/rose0/Projects/deepity/include/deepity/networks/DiscriminativePCNetwork.h:5, + from /home/rose0/Projects/deepity/include/deepity/utils/ModelIO.h:3, + from /home/rose0/Projects/deepity/src/ModelIO.cpp:1: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[146/152] Linking CXX static library bin/libDeepity.a +[147/152] Building CXX object CMakeFiles/MatMulLargeAsymmetricVerify.dir/tests/tMatMulLargeAsymmetricVerify.cpp.o +In file included from /home/rose0/Projects/deepity/include/deepity/backend/IComputeBackend.h:3, + from /home/rose0/Projects/deepity/include/deepity/backend/Backend.h:3, + from /home/rose0/Projects/deepity/tests/tMatMulLargeAsymmetricVerify.cpp:15: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[148/152] Building C object _deps/sleef-build/src/libm/CMakeFiles/sleefgnuabiavx512fsp.dir/sleefsimdsp.c.o +[149/152] Linking C static library _deps/sleef-build/lib/libsleefgnuabi.a +[150/152] Linking CXX executable bin/MatMulLargeAsymmetricVerify +[151/152] Building CXX object CMakeFiles/pydeepity.dir/bindings/pybinding.cpp.o +In file included from /home/rose0/Projects/deepity/bindings/pybinding.cpp:14: +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:731:24: warning: declaration of ‘void Deep::dGeluInto(float*, const float*, size_t)’ has a different exception specifier + 731 | static inline void dGeluInto(float *RESTRICT dst, const float *RESTRICT src, const size_t n) + | ^~~~~~~~~ +/home/rose0/Projects/deepity/include/deepity/utils/Activations.h:81:24: note: from previous declaration ‘void Deep::dGeluInto(float*, const float*, size_t) noexcept’ + 81 | static inline void dGeluInto(float *RESTRICT, const float *RESTRICT, size_t) noexcept; + | ^~~~~~~~~ +[152/152] Linking CXX shared module bin/pydeepity.cpython-314-x86_64-linux-gnu.so + + +=== Tests === +Test project /home/rose0/Projects/deepity/build/Release +No tests were found!!! diff --git a/mnist.py b/mnist.py index 879160d..6464d3b 100644 --- a/mnist.py +++ b/mnist.py @@ -39,26 +39,6 @@ def load_full_mnist(): Y_train[np.arange(y_train_labels.shape[0]), y_train_labels] = 1.0 - eps return X_train, Y_train, X_test, y_test_labels -def train_step_dfa(net, X, Y, inference_steps): - net.reset_state() - net.clamp_input(X) - net.project_forward() - net.get_terminal_layer().clamp_state(Y) - - net.calculate_terminal_error() - net.direct_feedback_update() - - for _ in range(inference_steps): - net.step() - - energy = 0.0 - for layer in net.layers: - energy += layer.calculate_state() - - net.update_weights() - net.get_terminal_layer().unclamp_state() - return energy - def main() -> None: SEED = int(sys.argv[1]) if len(sys.argv) > 1 else 7 EPOCHS = int(sys.argv[2]) if len(sys.argv) > 2 else 50 @@ -75,7 +55,7 @@ def main() -> None: DECAY_RATE = 0.94 print(f"\nBuilding network (784->512->512->10), seed={SEED}...") - net = DKPPCN(batch_size=BATCH_SIZE) + net = DKPPCN(batch_size=BATCH_SIZE, device="cpu") net.add_layer(784, 512, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="linear") # net.add_layer(512, 512, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="sigmoid") net.add_layer(512, TERMINAL_SIZE, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="sigmoid") @@ -112,7 +92,7 @@ def main() -> None: X_batch = X_shuf[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] Y_batch = Y_shuf[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] - energy = train_step_dfa(net, X_batch, Y_batch, INFERENCE_STEPS) + energy = net.train_step(X_batch, Y_batch, INFERENCE_STEPS) epoch_energy += energy N_ACC_BATCHES = 10 diff --git a/pydeepity/DKPPCN.py b/pydeepity/DKPPCN.py index b0cc519..e4b8845 100644 --- a/pydeepity/DKPPCN.py +++ b/pydeepity/DKPPCN.py @@ -37,8 +37,8 @@ class DKPPCN(dy.DirectKPPCNetwork): until that verification is done. """ - def __init__(self, batch_size: int) -> None: - super().__init__(batch_size) + def __init__(self, batch_size: int, device: str = "cpu") -> None: + super().__init__(batch_size, device) def add_layer( self, diff --git a/pydeepity/SimplePCN.py b/pydeepity/SimplePCN.py index 9a4c9e4..b23d409 100644 --- a/pydeepity/SimplePCN.py +++ b/pydeepity/SimplePCN.py @@ -19,9 +19,9 @@ class SimplePCN(dy.SimplePCNetwork): """ A Sequential Predictive Coding Network built from precision-stripped SimplePCLayers. """ - def __init__(self, batch_size: Optional[int] = None) -> None: + def __init__(self, batch_size: Optional[int] = None, device: str = "cpu") -> None: bsz = dy.auto_batch_size() if batch_size is None else batch_size - super().__init__(bsz) + super().__init__(bsz, device) def add_layer( self, @@ -44,8 +44,8 @@ def set_learning_rate(self, lr: float) -> None: def compile(self) -> None: super().compile() - def randomize_weights(self) -> None: - super().randomize_weights() + def randomize_weights(self, dist: str = "") -> None: + super().randomize_weights(dist) def train_step(self, X: npt.NDArray[np.float32], Y: npt.NDArray[np.float32], steps: int) -> float: self.reset_state() diff --git a/pydeepity/__pycache__/ConvolutionalPCN.cpython-312.pyc b/pydeepity/__pycache__/ConvolutionalPCN.cpython-312.pyc index 29095be..0af672c 100644 Binary files a/pydeepity/__pycache__/ConvolutionalPCN.cpython-312.pyc and b/pydeepity/__pycache__/ConvolutionalPCN.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/DKPPCN.cpython-312.pyc b/pydeepity/__pycache__/DKPPCN.cpython-312.pyc index 9ae2834..91e5e60 100644 Binary files a/pydeepity/__pycache__/DKPPCN.cpython-312.pyc and b/pydeepity/__pycache__/DKPPCN.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/DKPPCN.cpython-314.pyc b/pydeepity/__pycache__/DKPPCN.cpython-314.pyc index c94601f..3139c65 100644 Binary files a/pydeepity/__pycache__/DKPPCN.cpython-314.pyc and b/pydeepity/__pycache__/DKPPCN.cpython-314.pyc differ diff --git a/pydeepity/__pycache__/GaussSeidelPCN.cpython-312.pyc b/pydeepity/__pycache__/GaussSeidelPCN.cpython-312.pyc index 23a603f..52bb8b3 100644 Binary files a/pydeepity/__pycache__/GaussSeidelPCN.cpython-312.pyc and b/pydeepity/__pycache__/GaussSeidelPCN.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/SequentialPCN.cpython-312.pyc b/pydeepity/__pycache__/SequentialPCN.cpython-312.pyc index 4e46af2..23153d4 100644 Binary files a/pydeepity/__pycache__/SequentialPCN.cpython-312.pyc and b/pydeepity/__pycache__/SequentialPCN.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/SequentialPCN.cpython-314.pyc b/pydeepity/__pycache__/SequentialPCN.cpython-314.pyc index 76ca756..2a8661e 100644 Binary files a/pydeepity/__pycache__/SequentialPCN.cpython-314.pyc and b/pydeepity/__pycache__/SequentialPCN.cpython-314.pyc differ diff --git a/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-312.pyc b/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-312.pyc index 9eb0164..031033b 100644 Binary files a/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-312.pyc and b/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-314.pyc b/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-314.pyc index 9e0a8ad..cb111c4 100644 Binary files a/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-314.pyc and b/pydeepity/__pycache__/SimpleConvolutionalPCN.cpython-314.pyc differ diff --git a/pydeepity/__pycache__/SimplePCN.cpython-312.pyc b/pydeepity/__pycache__/SimplePCN.cpython-312.pyc index 86572cd..6872bb1 100644 Binary files a/pydeepity/__pycache__/SimplePCN.cpython-312.pyc and b/pydeepity/__pycache__/SimplePCN.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/SimplePCN.cpython-314.pyc b/pydeepity/__pycache__/SimplePCN.cpython-314.pyc index d0a1507..3408e8a 100644 Binary files a/pydeepity/__pycache__/SimplePCN.cpython-314.pyc and b/pydeepity/__pycache__/SimplePCN.cpython-314.pyc differ diff --git a/pydeepity/__pycache__/__init__.cpython-312.pyc b/pydeepity/__pycache__/__init__.cpython-312.pyc index d0ec1ba..e13c1b2 100644 Binary files a/pydeepity/__pycache__/__init__.cpython-312.pyc and b/pydeepity/__pycache__/__init__.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/_backend.cpython-312.pyc b/pydeepity/__pycache__/_backend.cpython-312.pyc index 824fb48..b1a8b1b 100644 Binary files a/pydeepity/__pycache__/_backend.cpython-312.pyc and b/pydeepity/__pycache__/_backend.cpython-312.pyc differ diff --git a/pydeepity/__pycache__/_backend.cpython-314.pyc b/pydeepity/__pycache__/_backend.cpython-314.pyc index e582871..94eab55 100644 Binary files a/pydeepity/__pycache__/_backend.cpython-314.pyc and b/pydeepity/__pycache__/_backend.cpython-314.pyc differ diff --git a/pydeepity/__pycache__/utils.cpython-312.pyc b/pydeepity/__pycache__/utils.cpython-312.pyc index 1fb7302..8cdb104 100644 Binary files a/pydeepity/__pycache__/utils.cpython-312.pyc and b/pydeepity/__pycache__/utils.cpython-312.pyc differ diff --git a/pydeepity/_backend.py b/pydeepity/_backend.py index 953875a..33f2bdc 100644 --- a/pydeepity/_backend.py +++ b/pydeepity/_backend.py @@ -1,6 +1,6 @@ try: - from . import pydeepity as dy + from . import pydeepity as dy except ImportError as e: - raise ImportError( - "Could not load the compiled Deepity C++ backend.\nEnsure the package was installed correctly or compiled for your architecture." - ) from e + raise ImportError( + "Could not load the compiled Deepity C++ backend.\nEnsure the package was installed correctly or compiled for your architecture." + ) from e diff --git a/src/DirectKPPCLayer.cpp b/src/DirectKPPCLayer.cpp index 670ca5a..eef2e54 100644 --- a/src/DirectKPPCLayer.cpp +++ b/src/DirectKPPCLayer.cpp @@ -1,11 +1,40 @@ #include "deepity/layers/DirectKPPCLayer.h" +#include "deepity/backend/CPUBackend.h" #include "deepity/utils/Optimize.h" +#include +#include namespace Deep { + namespace + { + ActivationType ToDerivativeType(ActivationType fwd) + { + switch (fwd) + { + case ActivationType::RELU: + return ActivationType::dRELU; + case ActivationType::SIGMOID: + return ActivationType::dSIGMOID; + case ActivationType::eSIGMOID: + return ActivationType::d_eSIGMOID; + case ActivationType::TANH: + return ActivationType::dTANH; + case ActivationType::LINEAR: + return ActivationType::dLINEAR; + default: + return ActivationType::NONE; + } + } + + void DeleteBackend(IComputeBackend *p) { delete p; } + void NoOpDeleter(IComputeBackend *) {} + } + DirectKPPCLayer::DirectKPPCLayer(size_t size, size_t nextSize, size_t terminalSize, size_t batchSize, float learningRate, float inferenceRate, float feedback, float lmbda, - ActivationType aType, ActivationType dType) + ActivationType aType, ActivationType dType, + IComputeBackend *backend) : size(size), nextSize(nextSize), terminalSize(terminalSize), @@ -14,18 +43,34 @@ namespace Deep ir(inferenceRate), fl(feedback), lmbda(lmbda), - activationType(aType) + activationType(aType), + backend(backend ? backend : new CPUBackend(), + backend ? NoOpDeleter : DeleteBackend) { this->activation = To_Fn(aType); this->activationDerivative = To_dFn(dType); + this->activationDerivativeInto = To_dFn2(dType); localArena = std::make_unique(GetRequiredFloats()); BindMemory(*localArena); } - // --- Setup --- + void DirectKPPCLayer::SetLearningRate(float learningRate) noexcept + { + lr = learningRate; + if (lr_device) + backend->CopyFromHost(lr_device, &lr, 1); + } + + void DirectKPPCLayer::SetFeedbackRate(float feedbackRate) noexcept + { + fl = feedbackRate; + if (fl_device) + backend->CopyFromHost(fl_device, &fl, 1); + } - void DirectKPPCLayer::BindMemory(MemoryArena &arena) + template + void DirectKPPCLayer::BindMemory(ArenaT &arena) { size_t own_state_size = batchSize * size; size_t out_state_size = batchSize * nextSize; @@ -34,8 +79,8 @@ namespace Deep z = arena.AllocateFloats(own_state_size); e = arena.AllocateFloats(own_state_size); - std::memset(z, 0, own_state_size * sizeof(float)); - std::memset(e, 0, own_state_size * sizeof(float)); + backend->Zero(z, own_state_size); + backend->Zero(e, own_state_size); if (nextSize > 0) { @@ -51,14 +96,17 @@ namespace Deep zFDeriv = arena.AllocateFloats(own_state_size); feedbackScratch = arena.AllocateFloats(own_state_size); - std::memset(b, 0, nextSize * sizeof(float)); - std::memset(mu, 0, out_state_size * sizeof(float)); - std::memset(cachedMu, 0, out_state_size * sizeof(float)); - std::memset(proj, 0, out_state_size * sizeof(float)); - std::memset(Psi, 0, direct_size * sizeof(float)); - std::memset(zF, 0, own_state_size * sizeof(float)); - std::memset(zFDeriv, 0, own_state_size * sizeof(float)); - std::memset(feedbackScratch, 0, own_state_size * sizeof(float)); + biasGradScratch = arena.AllocateFloats(nextSize); + + backend->Zero(b, nextSize); + backend->Zero(mu, out_state_size); + backend->Zero(cachedMu, out_state_size); + backend->Zero(proj, out_state_size); + backend->Zero(Psi, direct_size); + backend->Zero(zF, own_state_size); + backend->Zero(zFDeriv, own_state_size); + backend->Zero(feedbackScratch, own_state_size); + backend->Zero(biasGradScratch, nextSize); if (opt == OptimizerType::ADAM || opt == OptimizerType::ADAMW) { @@ -70,25 +118,44 @@ namespace Deep m_b = arena.AllocateFloats(nextSize); v_b = arena.AllocateFloats(nextSize); + backend->Zero(m_W, w_size); + backend->Zero(v_W, w_size); + backend->Zero(m_b, nextSize); + backend->Zero(v_b, nextSize); + backend->Zero(grad_W, w_size); + backend->Zero(grad_b, nextSize); + + t_device = reinterpret_cast(arena.AllocateFloats(1)); + lr_device = arena.AllocateFloats(1); + int zero = 0; + backend->CopyFromHost(reinterpret_cast(t_device), reinterpret_cast(&zero), 1); + backend->CopyFromHost(lr_device, &lr, 1); + } + + if (optPsi == OptimizerType::ADAM || optPsi == OptimizerType::ADAMW) + { grad_Psi = arena.AllocateFloats(direct_size); m_Psi = arena.AllocateFloats(direct_size); v_Psi = arena.AllocateFloats(direct_size); - std::memset(m_W, 0, w_size * sizeof(float)); - std::memset(v_W, 0, w_size * sizeof(float)); - std::memset(m_Psi, 0, direct_size * sizeof(float)); - - std::memset(m_b, 0, nextSize * sizeof(float)); - std::memset(v_b, 0, nextSize * sizeof(float)); - std::memset(v_Psi, 0, direct_size * sizeof(float)); + backend->Zero(m_Psi, direct_size); + backend->Zero(v_Psi, direct_size); + backend->Zero(grad_Psi, direct_size); - std::memset(grad_W, 0, w_size * sizeof(float)); - std::memset(grad_b, 0, nextSize * sizeof(float)); - std::memset(grad_Psi, 0, direct_size * sizeof(float)); + tPsi_device = reinterpret_cast(arena.AllocateFloats(1)); + fl_device = arena.AllocateFloats(1); + int zeroPsi = 0; + backend->CopyFromHost(reinterpret_cast(tPsi_device), reinterpret_cast(&zeroPsi), 1); + backend->CopyFromHost(fl_device, &fl, 1); } } - if (localArena && localArena.get() != &arena) + if constexpr (std::is_same_v) + { + if (localArena && localArena.get() != &arena) + localArena.reset(); + } + else { localArena.reset(); } @@ -103,7 +170,6 @@ namespace Deep size_t own_state_size = batchSize * size; size_t direct_size = terminalSize * size; - // dz_dt removed: state size drops from 3 arrays to 2 total += pad16(own_state_size) * 2; if (nextSize > 0) @@ -114,17 +180,22 @@ namespace Deep total += pad16(w_size); total += pad16(nextSize); total += pad16(out_state_size) * 3; - // E, prevZ, and bottom_up removed - total += pad16(own_state_size) * 3; // zF, zFDeriv, feedbackScratch + total += pad16(own_state_size) * 3; + total += pad16(direct_size); + total += pad16(nextSize); if (opt == OptimizerType::ADAM || opt == OptimizerType::ADAMW) { total += pad16(w_size) * 3; total += pad16(nextSize) * 3; - total += pad16(direct_size) * 3; + total += pad16(1) * 2; } - total += pad16(direct_size); + if (optPsi == OptimizerType::ADAM || optPsi == OptimizerType::ADAMW) + { + total += pad16(direct_size) * 3; + total += pad16(1) * 2; + } } return total; @@ -137,108 +208,37 @@ namespace Deep std::uniform_int_distribution seedDist; size_t Wsz = size * nextSize; + size_t Psisz = terminalSize * size; float limit = std::sqrt(2.0f / (size + nextSize)); float limPsi = std::sqrt(2.0f / (size + terminalSize)); - std::vector seeds(omp_get_max_threads()); - for (auto &s : seeds) - s = seedDist(seedGenerator); - -#pragma omp parallel if (!omp_in_parallel()) - { - std::mt19937 rng(seeds[omp_get_thread_num()]); - std::normal_distribution dist1(0.0f, limit); - std::normal_distribution dist2(0.0f, limPsi); - -#pragma omp for - for (ptrdiff_t i = 0; i < (ptrdiff_t)Wsz; ++i) - W[i] = dist1(rng); + uint32_t seedW = seedDist(seedGenerator); + uint32_t seedPsi = seedDist(seedGenerator); -#pragma omp for - for (ptrdiff_t i = 0; i < (ptrdiff_t)(terminalSize * size); ++i) - Psi[i] = dist2(rng); - } + backend->RandomizeNormal(W, Wsz, 0.0f, limit, seedW); + backend->RandomizeNormal(Psi, Psisz, 0.0f, limPsi, seedPsi); } - // --- Core DKP-PC Mechanics --- - - float DirectKPPCLayer::CalculateState() noexcept + float DirectKPPCLayer::CalculateState(bool needEnergy) noexcept { const size_t N = batchSize * size; if (layerBelow == nullptr) { - std::memset(e, 0, N * sizeof(float)); + backend->Zero(e, N); if (nextSize > 0) - { ComputeMuOnly(); - } return 0.0f; } - cblas_scopy(N, z, 1, e, 1); - cblas_saxpy(N, -1.0f, layerBelow->mu, 1, e, 1); - float totalEnergy = 0.0f; + if (needEnergy) + totalEnergy = backend->ComputeErrorAndEnergy(e, z, layerBelow->mu, N); + else + backend->ComputeError(e, z, layerBelow->mu, N); -#pragma omp parallel for schedule(static) reduction(+ : totalEnergy) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - const size_t offset = (size_t)batch * size; - size_t i = 0; - -#if defined(__AVX512F__) - __m512 half = _mm512_set1_ps(0.5f); - __m512 energy = _mm512_setzero_ps(); - size_t r = size % 16; - size_t simd_end = size - r; - for (; i < simd_end; i += 16) - { - __m512 e512 = _mm512_loadu_ps(&e[offset + i]); - energy = _mm512_fmadd_ps(half, _mm512_mul_ps(e512, e512), energy); - } - totalEnergy += _mm512_reduce_add_ps(energy); -#elif defined(__AVX2__) || defined(__AVX__) - __m256 half = _mm256_set1_ps(0.5f); - __m256 energy = _mm256_setzero_ps(); - size_t r = size % 8; - size_t simd_end = size - r; - for (; i < simd_end; i += 8) - { - __m256 e256 = _mm256_loadu_ps(&e[offset + i]); -#ifdef __FMA__ - energy = _mm256_fmadd_ps(half, _mm256_mul_ps(e256, e256), energy); -#else - energy = _mm256_add_ps(energy, _mm256_mul_ps(half, _mm256_mul_ps(e256, e256))); -#endif - } - totalEnergy += hsum256_ps(energy); -#elif defined(__SSE__) || defined(_M_AMD64) || defined(_M_X64) - __m128 half = _mm_set1_ps(0.5f); - __m128 energy = _mm_setzero_ps(); - size_t r = size % 4; - size_t simd_end = size - r; - for (; i < simd_end; i += 4) - { - __m128 e128 = _mm_loadu_ps(&e[offset + i]); -#ifdef __FMA__ - energy = _mm_fmadd_ps(half, _mm_mul_ps(e128, e128), energy); -#else - energy = _mm_add_ps(energy, _mm_mul_ps(half, _mm_mul_ps(e128, e128))); -#endif - } - totalEnergy += hsum128_ps(energy); -#endif - for (; i < size; ++i) - { - float err = e[offset + i]; - totalEnergy += 0.5f * err * err; - } - } if (nextSize > 0) - { ComputeMuOnly(); - } return totalEnergy; } @@ -253,47 +253,22 @@ namespace Deep if (isClamped && muCacheValid) { - cblas_scopy((int)Nout, cachedMu, 1, mu, 1); + backend->Copy(mu, cachedMu, Nout); return; } - cblas_scopy((int)N, z, 1, zF, 1); - switch (activationType) - { - case ActivationType::RELU: - Deep::relu(zF, N); - break; - case ActivationType::SIGMOID: - Deep::sigmoid(zF, N); - break; - case ActivationType::eSIGMOID: - Deep::e_sigmoid(zF, N); - break; - case ActivationType::TANH: - Deep::tanh(zF, N); - break; - case ActivationType::LINEAR: - Deep::linear(zF, N); - break; - default: - activation(zF, N); - break; // custom/unrecognized function pointer - } + backend->ActivationInto(activationType, zF, z, N); - cblas_sgemm( - CblasRowMajor, CblasNoTrans, CblasTrans, - batchSize, nextSize, size, - 1.0f, zF, size, W, size, 0.0f, mu, nextSize); + backend->MatMul( + /*transA=*/false, /*transB=*/true, + (int)batchSize, (int)nextSize, (int)size, + 1.0f, zF, (int)size, W, (int)size, 0.0f, mu, (int)nextSize); -#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - cblas_saxpy(nextSize, 1.0f, b, 1, mu + batch * nextSize, 1); - } + backend->AddBiasBroadcast(mu, b, batchSize, nextSize); if (isClamped) { - cblas_scopy((int)Nout, mu, 1, cachedMu, 1); + backend->Copy(cachedMu, mu, Nout); muCacheValid = true; } } @@ -309,61 +284,19 @@ namespace Deep { const float *e_above = layerAbove->GetErrors(); - switch (activationType) - { - case ActivationType::RELU: - Deep::dReluInto(zFDeriv, z, N); - break; - case ActivationType::SIGMOID: - Deep::dSigmoidInto(zFDeriv, z, N); - break; - case ActivationType::TANH: - Deep::dTanhInto(zFDeriv, z, N); - break; - case ActivationType::LINEAR: - Deep::dLinearInto(zFDeriv, z, N); - break; - default: - activationDerivativeInto(zFDeriv, z, N); - break; // custom/unrecognized, or eSIGMOID (no dedicated derivative variant exists) - } + backend->ActivationDerivativeInto(ToDerivativeType(activationType), zFDeriv, z, N); - cblas_sgemm( - CblasRowMajor, CblasNoTrans, CblasNoTrans, - batchSize, size, nextSize, - 1.0f, e_above, nextSize, W, size, - 0.0f, feedbackScratch, size); + backend->MatMul( + /*transA=*/false, /*transB=*/false, + (int)batchSize, (int)size, (int)nextSize, + 1.0f, e_above, (int)nextSize, W, (int)size, + 0.0f, feedbackScratch, (int)size); -#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - float *RESTRICT zPtr = z; - const float *RESTRICT feedbackPtr = feedbackScratch; - const float *RESTRICT zFDerivPtr = zFDeriv; - const float *RESTRICT ePtr = e; - size_t offset = (size_t)batch * size; - for (size_t i = 0; i < size; ++i) - { - size_t idx = offset + i; - // Fused dz_dt directly into the z update to save memory writes - zPtr[idx] += ir * ((feedbackPtr[idx] * zFDerivPtr[idx]) - ePtr[idx]); - } - } + backend->FusedStateUpdate(z, feedbackScratch, zFDeriv, e, N, ir); } - else // Output Layer + else { -#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - float *RESTRICT zPtr = z; - const float *RESTRICT ePtr = e; - size_t offset = (size_t)batch * size; - for (size_t i = 0; i < size; ++i) - { - size_t idx = offset + i; - zPtr[idx] += ir * -ePtr[idx]; - } - } + backend->AxpyInto(z, e, N, -ir); } } @@ -373,51 +306,49 @@ namespace Deep return; const float *local_grad = layerAbove->GetErrors(); - float grad_scale = -1.0f / batchSize; switch (opt) { case OptimizerType::SGD: { if (lmbda > 0.0f) - cblas_sscal((size_t)nextSize * size, 1.0f - lmbda, W, 1); + backend->Scale(W, (size_t)nextSize * size, 1.0f - lmbda); - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasNoTrans, - nextSize, size, batchSize, - lr / batchSize, local_grad, nextSize, zF, size, - 1.0f, W, size); + backend->MatMul( + /*transA=*/true, /*transB=*/false, + (int)nextSize, (int)size, (int)batchSize, + lr / batchSize, local_grad, (int)nextSize, zF, (int)size, + 1.0f, W, (int)size); float lr_batch = lr / batchSize; - for (int batch = 0; batch < batchSize; batch++) - cblas_saxpy(nextSize, lr_batch, local_grad + batch * nextSize, 1, b, 1); + backend->SumRows(biasGradScratch, local_grad, batchSize, nextSize); + backend->AxpyInto(b, biasGradScratch, nextSize, lr_batch); break; } case OptimizerType::ADAM: case OptimizerType::ADAMW: { - t++; + backend->IncrementCounter(t_device); size_t num_weights = (size_t)nextSize * size; float adam_scale = -1.0f; - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasNoTrans, - nextSize, size, batchSize, - adam_scale, local_grad, nextSize, zF, size, - 0.0f, grad_W, size); + backend->MatMul( + /*transA=*/true, /*transB=*/false, + (int)nextSize, (int)size, (int)batchSize, + adam_scale, local_grad, (int)nextSize, zF, (int)size, + 0.0f, grad_W, (int)size); - std::memset(grad_b, 0, nextSize * sizeof(float)); - for (int batch = 0; batch < batchSize; batch++) - cblas_saxpy(nextSize, adam_scale, local_grad + batch * nextSize, 1, grad_b, 1); + backend->SumRows(grad_b, local_grad, batchSize, nextSize); + backend->Scale(grad_b, nextSize, adam_scale); if (opt == OptimizerType::ADAMW) - Deep::AdamWUpdate(W, grad_W, m_W, v_W, num_weights, t, lr, lmbda); + backend->AdamWStep(W, grad_W, m_W, v_W, num_weights, t_device, lr_device, lmbda); else - Deep::AdamUpdate(W, grad_W, m_W, v_W, num_weights, t, lr); + backend->AdamStep(W, grad_W, m_W, v_W, num_weights, t_device, lr_device); - Deep::AdamUpdate(b, grad_b, m_b, v_b, nextSize, t, lr); + backend->AdamStep(b, grad_b, m_b, v_b, nextSize, t_device, lr_device); break; } } @@ -426,31 +357,31 @@ namespace Deep { case OptimizerType::SGD: { - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasNoTrans, - size, terminalSize, batchSize, - fl / batchSize, z, size, terminalLayer->GetErrors(), terminalSize, - 1.0f, Psi, terminalSize); + backend->MatMul( + /*transA=*/true, /*transB=*/false, + (int)size, (int)terminalSize, (int)batchSize, + fl / batchSize, z, (int)size, terminalLayer->GetErrors(), (int)terminalSize, + 1.0f, Psi, (int)terminalSize); break; } case OptimizerType::ADAMW: case OptimizerType::ADAM: { - tPsi++; + backend->IncrementCounter(tPsi_device); size_t num_weights_psi = size * terminalSize; float adam_scale = -1.0f; - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasNoTrans, - size, terminalSize, batchSize, - adam_scale, z, size, terminalLayer->GetErrors(), terminalSize, - 0.0f, grad_Psi, terminalSize); + backend->MatMul( + /*transA=*/true, /*transB=*/false, + (int)size, (int)terminalSize, (int)batchSize, + adam_scale, z, (int)size, terminalLayer->GetErrors(), (int)terminalSize, + 0.0f, grad_Psi, (int)terminalSize); if (optPsi == OptimizerType::ADAMW) - Deep::AdamWUpdate(Psi, grad_Psi, m_Psi, v_Psi, num_weights_psi, tPsi, fl, lmbda); + backend->AdamWStep(Psi, grad_Psi, m_Psi, v_Psi, num_weights_psi, tPsi_device, fl_device, lmbda); else - Deep::AdamUpdate(Psi, grad_Psi, m_Psi, v_Psi, num_weights_psi, tPsi, fl); + backend->AdamStep(Psi, grad_Psi, m_Psi, v_Psi, num_weights_psi, tPsi_device, fl_device); break; } @@ -459,33 +390,30 @@ namespace Deep void DirectKPPCLayer::DirectFeedbackUpdate() noexcept { - // If the layer above has no Psi weights (i.e. it is the terminal layer), skip DFA if (layerAbove == nullptr || layerAbove->GetDirectFeedbackWeights() == nullptr) return; // proj = terminalLayer->GetErrors() @ layerAbove->GetDirectFeedbackWeights()^T - cblas_sgemm( - CblasRowMajor, CblasNoTrans, CblasTrans, - batchSize, nextSize, terminalSize, - 1.0f, terminalLayer->GetErrors(), terminalSize, - layerAbove->GetDirectFeedbackWeights(), terminalSize, - 0.0f, proj, nextSize); - - // W += fl * proj^T @ zF - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasTrans, - nextSize, size, batchSize, - fl / batchSize, proj, nextSize, - zF, size, - 1.0f, W, size); + backend->MatMul( + /*transA=*/false, /*transB=*/true, + (int)batchSize, (int)nextSize, (int)terminalSize, + 1.0f, terminalLayer->GetErrors(), (int)terminalSize, + layerAbove->GetDirectFeedbackWeights(), (int)terminalSize, + 0.0f, proj, (int)nextSize); + + // W += fl/batchSize * proj^T @ zF + backend->MatMul( + /*transA=*/true, /*transB=*/false, + (int)nextSize, (int)size, (int)batchSize, + fl / batchSize, proj, (int)nextSize, + zF, (int)size, + 1.0f, W, (int)size); } - // --- Getters / Setters --- - void DirectKPPCLayer::ClampState(const std::vector &inputData) noexcept { - size_t copySize = (std::min)(inputData.size(), (size_t)(batchSize * size)) * sizeof(float); - memcpy(z, inputData.data(), copySize); + size_t copyFloats = (std::min)(inputData.size(), (size_t)(batchSize * size)); + backend->CopyFromHost(z, inputData.data(), copyFloats); isClamped = true; muCacheValid = false; } @@ -498,6 +426,11 @@ namespace Deep void DirectKPPCLayer::ResetState() noexcept { size_t N = (size_t)batchSize * size; - std::memset(z, 0, N * sizeof(float)); + backend->Zero(z, N); } -} + + template void DirectKPPCLayer::BindMemory(MemoryArena &arena); +#if defined(DEEPITY_USE_CUDA) + template void DirectKPPCLayer::BindMemory(DeviceMemoryArena &arena); +#endif +} \ No newline at end of file diff --git a/src/DirectKPPCNetwork.cpp b/src/DirectKPPCNetwork.cpp index e9c7db7..c4ea5ca 100644 --- a/src/DirectKPPCNetwork.cpp +++ b/src/DirectKPPCNetwork.cpp @@ -1,20 +1,22 @@ #include -#ifdef DEEPITY_USE_MKL -#include -#else -#include -#endif +#include +#include +#include namespace Deep { - DirectKPPCNetwork::DirectKPPCNetwork(int batchSize) noexcept - : batchSize(batchSize) {} + DirectKPPCNetwork::DirectKPPCNetwork(int batchSize, DeviceType device) noexcept + : device(device), batchSize(batchSize) + { + backend = CreateBackend(device); + } void DirectKPPCNetwork::AddLayer(size_t size, size_t nextSize, size_t terminalSize, float lr, float ir, float fl, float lmbda, ActivationType aType, ActivationType dType) { - std::unique_ptr l = std::make_unique(size, nextSize, terminalSize, batchSize, lr, ir, fl, lmbda, aType, dType); + std::unique_ptr l = std::make_unique( + size, nextSize, terminalSize, batchSize, lr, ir, fl, lmbda, aType, dType, backend.get()); if (!layers.empty()) { @@ -48,33 +50,43 @@ namespace Deep { layers[i]->ComputeMuOnly(); + // Skip a layer that's already clamped (the terminal layer, + // once ClampState(y) has run) -- overwriting its real target + // with a forward-projected guess is exactly the bug + // SimplePCNetwork::ProjectForward() had, fixed here before + // it gets exercised for the first time by moving this call + // inside graph capture, which requires ClampState(y) to run + // BEFORE this, not after, for the capture ordering to work. + if (layers[i + 1]->IsClamped()) + continue; + const float *mu = layers[i]->GetMu(); float *nextZ = layers[i + 1]->GetBeliefs(); size_t n = layers[i]->GetBatchSize() * layers[i]->GetOutputSize(); - std::memcpy(nextZ, mu, n * sizeof(float)); + backend->Copy(nextZ, mu, n); } } - float DirectKPPCNetwork::CalculateTerminalError() noexcept - { - return GetTerminalLayer()->CalculateState(); - } - void DirectKPPCNetwork::DirectFeedbackUpdate() noexcept { for (size_t i = 0; i < layers.size() - 1; ++i) layers[i]->DirectFeedbackUpdate(); } - float DirectKPPCNetwork::Step() noexcept + float DirectKPPCNetwork::CalculateTerminalError() noexcept + { + return GetTerminalLayer()->CalculateState(false); + } + + float DirectKPPCNetwork::Step(bool needEnergy) noexcept { float e = 0.0f; for (auto &l : layers) - e += l->CalculateState(); + e += l->CalculateState(needEnergy); for (auto &l : layers) l->UpdateState(); - return e; + return needEnergy ? e : 0.0f; } void DirectKPPCNetwork::UpdateWeights() noexcept @@ -89,28 +101,65 @@ namespace Deep { ResetState(); Clamp(x); - ProjectForward(); + // Moved BEFORE ProjectForward() -- required for ProjectForward's + // new IsClamped() guard to actually protect the terminal layer, + // and required so ProjectForward() can safely move inside the + // captured region below. GetTerminalLayer()->ClampState(y); - CalculateTerminalError(); - DirectFeedbackUpdate(); - - for (int t = 0; t < inferenceSteps; t++) - Step(); + if (device == DeviceType::DEVICE_GPU) + { + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); // now inside capture -- see note above + CalculateTerminalError(); + DirectFeedbackUpdate(); + 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(); + } + else + { + ProjectForward(); + CalculateTerminalError(); + DirectFeedbackUpdate(); + for (int t = 0; t < inferenceSteps; t++) + Step(false); + UpdateWeights(); + } + } + else + { + ProjectForward(); + CalculateTerminalError(); + DirectFeedbackUpdate(); + for (int t = 0; t < inferenceSteps; t++) + Step(false); + UpdateWeights(); + } - // Sync e/mu/zF to the TRUE final z before UpdateWeights() reads them. - // Step()'s own CalculateState() each iteration reflects z BEFORE that - // iteration's UpdateState() moves it, so after the loop exits, e/mu/zF - // are one iteration stale relative to the final z -- exactly the bug - // the test's added TotalEnergy() call worked around. This also gives a - // more accurate finalEnergy for free, computed at the true final state. float finalEnergy = 0.0f; for (auto &l : layers) - finalEnergy += l->CalculateState(); + finalEnergy += l->CalculateState(true); - UpdateWeights(); GetTerminalLayer()->UnclampState(); - return finalEnergy; } @@ -130,13 +179,15 @@ namespace Deep const float *beliefs = terminal->GetBeliefs(); size_t count = terminal->GetBatchSize() * terminal->GetInputSize(); - return std::vector(beliefs, beliefs + count); + std::vector result(count); + backend->CopyToHost(result.data(), beliefs, count); + return result; } void DirectKPPCNetwork::Compile() { #pragma omp parallel - { // Broadcast FTZ/DAZ hardware flags to ALL OpenMP worker threads + { _MM_SET_FLUSH_ZERO_MODE(_MM_FLUSH_ZERO_ON); _MM_SET_DENORMALS_ZERO_MODE(_MM_DENORMALS_ZERO_ON); } @@ -144,12 +195,26 @@ namespace Deep for (auto &layer : layers) total_floats_needed += layer->GetRequiredFloats(); - arena = std::make_unique(total_floats_needed, false); // TODO: Consider adding huge pages as a param - - for (auto &layer : layers) + if (device == DeviceType::DEVICE_CPU) { - layer->BindMemory(*arena); - layer->SetTerminalLayer(layers.back().get()); + cpuArena = std::make_unique(total_floats_needed, false); + for (auto &layer : layers) + { + layer->BindMemory(*cpuArena); + layer->SetTerminalLayer(layers.back().get()); + } } +#if defined(DEEPITY_USE_CUDA) + else + { + backend->PrepareForBatchSize(batchSize); + gpuArena = std::make_unique(backend.get(), total_floats_needed); + for (auto &layer : layers) + { + layer->BindMemory(*gpuArena); + layer->SetTerminalLayer(layers.back().get()); + } + } +#endif } -} +} \ No newline at end of file diff --git a/src/SimplePCLayer.cpp b/src/SimplePCLayer.cpp index 8f8d798..0d6f462 100644 --- a/src/SimplePCLayer.cpp +++ b/src/SimplePCLayer.cpp @@ -1,26 +1,48 @@ #include +#include #include #include -#include #include -#ifdef DEEPITY_USE_MKL -#include -#else -#include -#endif #include -#include #include -#include #include +#include namespace Deep { + namespace + { + ActivationType ToDerivativeType(ActivationType fwd) + { + switch (fwd) + { + case ActivationType::RELU: + return ActivationType::dRELU; + case ActivationType::SIGMOID: + return ActivationType::dSIGMOID; + case ActivationType::eSIGMOID: + return ActivationType::d_eSIGMOID; + case ActivationType::TANH: + return ActivationType::dTANH; + case ActivationType::LINEAR: + return ActivationType::dLINEAR; + default: + return ActivationType::NONE; + } + } + + void DeleteBackend(IComputeBackend *p) { delete p; } + void NoOpDeleter(IComputeBackend *) {} + } + SimplePCLayer::SimplePCLayer(size_t size, size_t nextSize, size_t batchSize, float learningRate, float inferenceRate, float lmbda, void (*act)(float *, size_t), - void (*dAct)(float *, size_t, bool)) - : batchSize(batchSize), lr(learningRate), ir(inferenceRate), lmbda(lmbda), isClamped(false), + void (*dAct)(float *, size_t, bool), + IComputeBackend *backend) + : backend(backend ? backend : new CPUBackend(), + backend ? NoOpDeleter : DeleteBackend), + batchSize(batchSize), lr(learningRate), ir(inferenceRate), lmbda(lmbda), isClamped(false), layerAbove(nullptr), layerBelow(nullptr), activation(act), activationDerivative(dAct), activationType(To_AType(act)), opt(OptimizerType::SGD) { this->size = size; @@ -34,8 +56,11 @@ namespace Deep SimplePCLayer::SimplePCLayer(size_t size, size_t nextSize, size_t batchSize, float learningRate, float inferenceRate, float lmbda, - ActivationType aType, ActivationType dType) - : batchSize(batchSize), lr(learningRate), ir(inferenceRate), lmbda(lmbda), isClamped(false), + ActivationType aType, ActivationType dType, + IComputeBackend *backend) + : backend(backend ? backend : new CPUBackend(), + backend ? NoOpDeleter : DeleteBackend), + batchSize(batchSize), lr(learningRate), ir(inferenceRate), lmbda(lmbda), isClamped(false), layerAbove(nullptr), layerBelow(nullptr), activationType(aType) { this->activation = To_Fn(aType); @@ -49,104 +74,54 @@ namespace Deep BindMemory(*localArena); } - void SimplePCLayer::RandomizeWeights(std::mt19937 &seedGenerator) noexcept + void SimplePCLayer::RandomizeWeights(std::mt19937 &seedGenerator, const char *distribution) noexcept { - std::uniform_int_distribution seedDist; - size_t Wsz = size * nextSize; - float limit = std::sqrt(2.0f / (size + nextSize)); - - std::vector seeds(omp_get_max_threads()); - for (auto &s : seeds) - s = seedDist(seedGenerator); + char name[16] = {0}; + float a = 0.0f, b = 1.0f; -#pragma omp parallel if (!omp_in_parallel()) + if (sscanf(distribution, "%15[^(](%f,%f)", name, &a, &b) != 3) { - std::mt19937 rng(seeds[omp_get_thread_num()]); - std::normal_distribution dist(0.0f, limit); - -#pragma omp for - for (ptrdiff_t i = 0; i < (ptrdiff_t)Wsz; ++i) - W[i] = dist(rng); + // Malformed string -- fall back to this class's original, + // validated default (He/Xavier-style normal init) rather than + // silently doing nothing. + name[0] = '\0'; + strcpy(name, "normal"); + a = 0.0f; + b = std::sqrt(2.0f / (size + nextSize)); } + + std::uniform_int_distribution seedDist; + size_t Wsz = size * nextSize; + uint32_t seed = seedDist(seedGenerator); + + if (strcmp(name, "normal") == 0) + backend->RandomizeNormal(W, Wsz, a, b, seed); + else if (strcmp(name, "uniform") == 0) + backend->RandomizeUniform(W, Wsz, a, b, seed); + else + backend->RandomizeNormal(W, Wsz, 0.0f, std::sqrt(2.0f / (size + nextSize)), seed); } - float SimplePCLayer::CalculateState() noexcept + float SimplePCLayer::CalculateState(bool needEnergy) noexcept { const size_t N = batchSize * size; if (layerBelow == nullptr) { - std::memset(e, 0, N * sizeof(float)); + backend->Zero(e, N); if (nextSize > 0) - { ComputeMuOnly(); - } return 0.0f; } - cblas_scopy(N, z, 1, e, 1); - cblas_saxpy(N, -1.0f, layerBelow->mu, 1, e, 1); - float totalEnergy = 0.0f; + if (needEnergy) + totalEnergy = backend->ComputeErrorAndEnergy(e, z, layerBelow->mu, N); + else + backend->ComputeError(e, z, layerBelow->mu, N); -#pragma omp parallel for schedule(static) reduction(+ : totalEnergy) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - const size_t offset = (size_t)batch * size; - size_t i = 0; - -#if defined(__AVX512F__) - __m512 half = _mm512_set1_ps(0.5f); - __m512 energy = _mm512_setzero_ps(); - size_t r = size % 16; - size_t simd_end = size - r; - for (; i < simd_end; i += 16) - { - __m512 e512 = _mm512_loadu_ps(&e[offset + i]); - energy = _mm512_fmadd_ps(half, _mm512_mul_ps(e512, e512), energy); - } - totalEnergy += _mm512_reduce_add_ps(energy); -#elif defined(__AVX2__) || defined(__AVX__) - __m256 half = _mm256_set1_ps(0.5f); - __m256 energy = _mm256_setzero_ps(); - size_t r = size % 8; - size_t simd_end = size - r; - for (; i < simd_end; i += 8) - { - __m256 e256 = _mm256_loadu_ps(&e[offset + i]); -#ifdef __FMA__ - energy = _mm256_fmadd_ps(half, _mm256_mul_ps(e256, e256), energy); -#else - energy = _mm256_add_ps(energy, _mm256_mul_ps(half, _mm256_mul_ps(e256, e256))); -#endif - } - totalEnergy += hsum256_ps(energy); -#elif defined(__SSE__) || defined(_M_AMD64) || defined(_M_X64) - __m128 half = _mm_set1_ps(0.5f); - __m128 energy = _mm_setzero_ps(); - size_t r = size % 4; - size_t simd_end = size - r; - for (; i < simd_end; i += 4) - { - __m128 e128 = _mm_loadu_ps(&e[offset + i]); -#ifdef __FMA__ - energy = _mm_fmadd_ps(half, _mm_mul_ps(e128, e128), energy); -#else - energy = _mm_add_ps(energy, _mm_mul_ps(half, _mm_mul_ps(e128, e128))); -#endif - } - totalEnergy += hsum128_ps(energy); -#endif - for (; i < size; ++i) - { - float err = e[offset + i]; - totalEnergy += 0.5f * err * err; - } - } if (nextSize > 0) - { ComputeMuOnly(); - } return totalEnergy; } @@ -161,47 +136,22 @@ namespace Deep if (isClamped && muCacheValid) { - cblas_scopy((int)Nout, cachedMu, 1, mu, 1); + backend->Copy(mu, cachedMu, Nout); return; } - cblas_scopy((int)N, z, 1, zF, 1); - switch (activationType) - { - case ActivationType::RELU: - Deep::relu(zF, N); - break; - case ActivationType::SIGMOID: - Deep::sigmoid(zF, N); - break; - case ActivationType::eSIGMOID: - Deep::e_sigmoid(zF, N); - break; - case ActivationType::TANH: - Deep::tanh(zF, N); - break; - case ActivationType::LINEAR: - Deep::linear(zF, N); - break; - default: - activation(zF, N); - break; // custom/unrecognized function pointer - } + backend->ActivationInto(activationType, zF, z, N); - cblas_sgemm( - CblasRowMajor, CblasNoTrans, CblasTrans, - batchSize, nextSize, size, - 1.0f, zF, size, W, size, 0.0f, mu, nextSize); + backend->MatMul( + false, true, + (int)batchSize, (int)nextSize, (int)size, + 1.0f, zF, (int)size, W, (int)size, 0.0f, mu, (int)nextSize); -#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - cblas_saxpy(nextSize, 1.0f, b, 1, mu + batch * nextSize, 1); - } + backend->AddBiasBroadcast(mu, b, batchSize, nextSize); if (isClamped) { - cblas_scopy((int)Nout, mu, 1, cachedMu, 1); + backend->Copy(cachedMu, mu, Nout); muCacheValid = true; } } @@ -217,61 +167,21 @@ namespace Deep { const float *e_above = layerAbove->GetErrors(); - switch (activationType) - { - case ActivationType::RELU: - Deep::dReluInto(zFDeriv, z, N); - break; - case ActivationType::SIGMOID: - Deep::dSigmoidInto(zFDeriv, z, N); - break; - case ActivationType::TANH: - Deep::dTanhInto(zFDeriv, z, N); - break; - case ActivationType::LINEAR: - Deep::dLinearInto(zFDeriv, z, N); - break; - default: - activationDerivativeInto(zFDeriv, z, N); - break; // custom/unrecognized, or eSIGMOID (no dedicated derivative variant exists) - } + backend->ActivationDerivativeInto(ToDerivativeType(activationType), zFDeriv, z, N); - cblas_sgemm( - CblasRowMajor, CblasNoTrans, CblasNoTrans, - batchSize, size, nextSize, - 1.0f, e_above, nextSize, W, size, - 0.0f, feedbackScratch, size); + backend->MatMul( + /*transA=*/false, /*transB=*/false, + (int)batchSize, (int)size, (int)nextSize, + 1.0f, e_above, (int)nextSize, W, (int)size, + 0.0f, feedbackScratch, (int)size); -#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - float *RESTRICT zPtr = z; - const float *RESTRICT feedbackPtr = feedbackScratch; - const float *RESTRICT zFDerivPtr = zFDeriv; - const float *RESTRICT ePtr = e; - size_t offset = (size_t)batch * size; - for (size_t i = 0; i < size; ++i) - { - size_t idx = offset + i; - // Fused dz_dt directly into the z update to save memory writes - zPtr[idx] += ir * ((feedbackPtr[idx] * zFDerivPtr[idx]) - ePtr[idx]); - } - } + backend->FusedStateUpdate(z, feedbackScratch, zFDeriv, e, N, ir); } else // Output Layer { -#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) - for (int batch = 0; batch < batchSize; ++batch) - { - float *RESTRICT zPtr = z; - const float *RESTRICT ePtr = e; - size_t offset = (size_t)batch * size; - for (size_t i = 0; i < size; ++i) - { - size_t idx = offset + i; - zPtr[idx] += ir * -ePtr[idx]; - } - } + // z[i] += ir * (-e[i]) == z += (-ir) * e, i.e. AxpyInto + // with alpha = -ir. + backend->AxpyInto(z, e, N, -ir); } } @@ -287,48 +197,44 @@ namespace Deep case OptimizerType::SGD: { if (lmbda > 0.0f) - cblas_sscal((size_t)nextSize * size, 1.0f - lmbda, W, 1); + backend->Scale(W, (size_t)nextSize * size, 1.0f - lmbda); - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasNoTrans, - nextSize, size, batchSize, - lr / batchSize, local_grad, nextSize, zF, size, - 1.0f, W, size); + backend->MatMul( + true, false, + (int)nextSize, (int)size, (int)batchSize, + lr / batchSize, local_grad, (int)nextSize, zF, (int)size, + 1.0f, W, (int)size); float lr_batch = lr / batchSize; for (int batch = 0; batch < batchSize; batch++) - cblas_saxpy(nextSize, lr_batch, local_grad + batch * nextSize, 1, b, 1); + backend->AxpyInto(b, local_grad + batch * nextSize, nextSize, lr_batch); break; } case OptimizerType::ADAM: case OptimizerType::ADAMW: { - t++; + backend->IncrementCounter(t_device); // was: t++ size_t num_weights = (size_t)nextSize * size; float grad_scale = -1.0f; - cblas_sgemm( - CblasRowMajor, CblasTrans, CblasNoTrans, - nextSize, size, batchSize, - grad_scale, local_grad, nextSize, zF, size, - 0.0f, grad_W, size); + backend->MatMul( + true, false, + (int)nextSize, (int)size, (int)batchSize, + grad_scale, local_grad, (int)nextSize, zF, (int)size, + 0.0f, grad_W, (int)size); - std::memset(grad_b, 0, nextSize * sizeof(float)); + backend->Zero(grad_b, nextSize); for (int batch = 0; batch < batchSize; batch++) - cblas_saxpy(nextSize, grad_scale, local_grad + batch * nextSize, 1, grad_b, 1); + backend->AxpyInto(grad_b, local_grad + batch * nextSize, nextSize, grad_scale); if (opt == OptimizerType::ADAMW) - { - Deep::AdamWUpdate(W, grad_W, m_W, v_W, num_weights, t, lr, lmbda); - } + backend->AdamWStep(W, grad_W, m_W, v_W, num_weights, t_device, lr_device, lmbda); else - { - Deep::AdamUpdate(W, grad_W, m_W, v_W, num_weights, t, lr); - } + backend->AdamStep(W, grad_W, m_W, v_W, num_weights, t_device, lr_device); - Deep::AdamUpdate(b, grad_b, m_b, v_b, nextSize, t, lr); + backend->AdamStep(b, grad_b, m_b, v_b, nextSize, t_device, lr_device); break; } } @@ -337,13 +243,13 @@ namespace Deep void SimplePCLayer::ResetState() noexcept { size_t N = (size_t)batchSize * size; - std::memset(z, 0, N * sizeof(float)); + backend->Zero(z, N); } void SimplePCLayer::ClampState(const std::vector &inputData) noexcept { - size_t copySize = (std::min)(inputData.size(), (size_t)(batchSize * size)) * sizeof(float); - memcpy(z, inputData.data(), copySize); + size_t copyFloats = (std::min)(inputData.size(), (size_t)(batchSize * size)); + backend->CopyFromHost(z, inputData.data(), copyFloats); isClamped = true; muCacheValid = false; } @@ -361,7 +267,6 @@ namespace Deep size_t total = 0; size_t own_state_size = (size_t)batchSize * size; - // dz_dt removed: state size drops from 3 arrays to 2 total += pad16(own_state_size) * 2; if (nextSize > 0) @@ -372,20 +277,27 @@ namespace Deep total += pad16(w_size); total += pad16(nextSize); total += pad16(out_state_size) * 2; - // E, prevZ, and bottom_up removed - total += pad16(own_state_size) * 3; // zF, zFDeriv, feedbackScratch + total += pad16(own_state_size) * 3; if (opt == OptimizerType::ADAM || opt == OptimizerType::ADAMW) { total += pad16(w_size) * 3; total += pad16(nextSize) * 3; + total += pad16(1) * 2; // t_device, lr_device } } return total; } - void SimplePCLayer::BindMemory(MemoryArena &arena) + void SimplePCLayer::SetLearningRate(float lr) noexcept + { + this->lr = lr; + backend->CopyFromHost(lr_device, &this->lr, 1); + } + + template + void SimplePCLayer::BindMemory(ArenaT &arena) { size_t own_state_size = (size_t)batchSize * size; size_t out_state_size = (size_t)batchSize * nextSize; @@ -393,8 +305,8 @@ namespace Deep z = arena.AllocateFloats(own_state_size); e = arena.AllocateFloats(own_state_size); - std::memset(z, 0, own_state_size * sizeof(float)); - std::memset(e, 0, own_state_size * sizeof(float)); + backend->Zero(z, own_state_size); + backend->Zero(e, own_state_size); if (nextSize > 0) { @@ -408,12 +320,12 @@ namespace Deep zFDeriv = arena.AllocateFloats(own_state_size); feedbackScratch = arena.AllocateFloats(own_state_size); - std::memset(b, 0, nextSize * sizeof(float)); - std::memset(mu, 0, out_state_size * sizeof(float)); - std::memset(cachedMu, 0, out_state_size * sizeof(float)); - std::memset(zF, 0, own_state_size * sizeof(float)); - std::memset(zFDeriv, 0, own_state_size * sizeof(float)); - std::memset(feedbackScratch, 0, own_state_size * sizeof(float)); + backend->Zero(b, nextSize); + backend->Zero(mu, out_state_size); + backend->Zero(cachedMu, out_state_size); + backend->Zero(zF, own_state_size); + backend->Zero(zFDeriv, own_state_size); + backend->Zero(feedbackScratch, own_state_size); if (opt == OptimizerType::ADAM || opt == OptimizerType::ADAMW) { @@ -425,19 +337,35 @@ namespace Deep m_b = arena.AllocateFloats(nextSize); v_b = arena.AllocateFloats(nextSize); - std::memset(m_W, 0, w_size * sizeof(float)); - std::memset(v_W, 0, w_size * sizeof(float)); - std::memset(m_b, 0, nextSize * sizeof(float)); - std::memset(v_b, 0, nextSize * sizeof(float)); + backend->Zero(m_W, w_size); + backend->Zero(v_W, w_size); + backend->Zero(m_b, nextSize); + backend->Zero(v_b, nextSize); + + backend->Zero(grad_W, w_size); + backend->Zero(grad_b, nextSize); - std::memset(grad_W, 0, w_size * sizeof(float)); - std::memset(grad_b, 0, nextSize * sizeof(float)); + t_device = reinterpret_cast(arena.AllocateFloats(1)); + lr_device = arena.AllocateFloats(1); + + int zero = 0; + backend->CopyFromHost(reinterpret_cast(t_device), reinterpret_cast(&zero), 1); + backend->CopyFromHost(lr_device, &lr, 1); } } - - if (localArena && localArena.get() != &arena) + if constexpr (std::is_same_v) + { + if (localArena && localArena.get() != &arena) + localArena.reset(); + } + else { localArena.reset(); } } -} + + template void SimplePCLayer::BindMemory(MemoryArena &arena); +#if defined(DEEPITY_USE_CUDA) + template void SimplePCLayer::BindMemory(DeviceMemoryArena &arena); +#endif +} \ No newline at end of file diff --git a/src/SimplePCNetwork.cpp b/src/SimplePCNetwork.cpp index 960c620..baebb83 100644 --- a/src/SimplePCNetwork.cpp +++ b/src/SimplePCNetwork.cpp @@ -2,15 +2,21 @@ #include #include #include +#include namespace Deep { - SimplePCNetwork::SimplePCNetwork(int batchSize) noexcept : batchSize(batchSize) {} + SimplePCNetwork::SimplePCNetwork(int batchSize, DeviceType device) noexcept + : device(device), batchSize(batchSize) + { + backend = CreateBackend(device); + } void SimplePCNetwork::AddLayer(int size, int nextSize, float lr, float ir, float lmbda, void (*act)(float *, size_t), void (*dAct)(float *, size_t, bool)) { - std::unique_ptr l = std::make_unique(size, nextSize, batchSize, lr, ir, lmbda, act, dAct); + std::unique_ptr l = std::make_unique( + size, nextSize, batchSize, lr, ir, lmbda, act, dAct, backend.get()); if (!layers.empty()) { @@ -23,7 +29,8 @@ namespace Deep void SimplePCNetwork::AddLayer(int size, int nextSize, float lr, float ir, float lmbda, ActivationType aType, ActivationType dType) { - std::unique_ptr l = std::make_unique(size, nextSize, batchSize, lr, ir, lmbda, aType, dType); + std::unique_ptr l = std::make_unique( + size, nextSize, batchSize, lr, ir, lmbda, aType, dType, backend.get()); if (!layers.empty()) { @@ -33,6 +40,12 @@ namespace Deep layers.push_back(std::move(l)); } + void SimplePCNetwork::RandomizeWeights(std::mt19937 &rng, const char *distribution) + { + for (auto &l : layers) + l->RandomizeWeights(rng, distribution); + } + void SimplePCNetwork::RandomizeWeights(std::mt19937 &rng) { for (auto &l : layers) @@ -50,12 +63,12 @@ namespace Deep layers.front()->ClampState(input); } - float SimplePCNetwork::CalculateState() + float SimplePCNetwork::CalculateState(bool needEnergy) { float e = 0.0f; for (size_t i = 0; i < layers.size(); i++) - e += layers[i]->CalculateState(); - return e; + e += layers[i]->CalculateState(needEnergy); + return needEnergy ? e : 0.0f; } void SimplePCNetwork::UpdateState() @@ -78,11 +91,11 @@ namespace Deep for (int t = 0; t < inferenceSteps; t++) { - CalculateState(); + CalculateState(false); UpdateState(); } - float finalEnergy = CalculateState(); + float finalEnergy = CalculateState(true); UpdateWeights(); GetTerminalLayer()->UnclampState(); @@ -105,7 +118,9 @@ namespace Deep const float *beliefs = terminal->GetBeliefs(); size_t count = terminal->GetBatchSize() * terminal->GetInputSize(); - return std::vector(beliefs, beliefs + count); + std::vector result(count); + backend->CopyToHost(result.data(), beliefs, count); + return result; } void SimplePCNetwork::ProjectForward() noexcept @@ -114,30 +129,74 @@ namespace Deep { layers[i]->ComputeMuOnly(); + if (layers[i + 1]->IsClamped()) + continue; // never overwrite a clamped layer's real target with a forward guess + const float *mu = layers[i]->GetMu(); float *nextZ = layers[i + 1]->GetBeliefs(); size_t n = layers[i]->GetBatchSize() * layers[i]->GetOutputSize(); - std::memcpy(nextZ, mu, n * sizeof(float)); + backend->Copy(nextZ, mu, n); } } - float SimplePCNetwork::TrainStepWithProjection(const std::vector &x, const std::vector &y, int inferenceSteps) + float SimplePCNetwork::TrainStepWithProjection(const std::vector &x, const std::vector &y, int inferenceSteps, bool computeEnergy) { ResetState(); Clamp(x); - ProjectForward(); GetTerminalLayer()->ClampState(y); - float finalEnergy = 0.0f; - for (int t = 0; t < inferenceSteps; ++t) + if (device == DeviceType::DEVICE_GPU) { - CalculateState(); - UpdateState(); + if (!graphCaptured || capturedInferenceSteps != inferenceSteps) + { + backend->BeginGraphCapture(); + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + bool captureOk = backend->EndGraphCapture(); + + if (captureOk) + { + graphCaptured = true; + capturedInferenceSteps = inferenceSteps; + } + else + { + std::cerr << "Graph capture failed -- falling back to non-graph execution for this call.\n"; + } + } + + if (graphCaptured) + { + backend->ReplayGraph(); + } + else + { + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); + } + } + else + { + ProjectForward(); + for (int t = 0; t < inferenceSteps; ++t) + { + CalculateState(false); + UpdateState(); + } + UpdateWeights(); } - finalEnergy = CalculateState(); - UpdateWeights(); + float finalEnergy = CalculateState(computeEnergy); GetTerminalLayer()->UnclampState(); return finalEnergy; @@ -147,11 +206,11 @@ namespace Deep { ResetState(); Clamp(x); - ProjectForward(); // Add your forward projection initialization here + ProjectForward(); for (int t = 0; t < inferenceSteps; t++) { - CalculateState(); + CalculateState(false); UpdateState(); } @@ -159,7 +218,9 @@ namespace Deep const float *beliefs = terminal->GetBeliefs(); size_t count = terminal->GetBatchSize() * terminal->GetInputSize(); - return std::vector(beliefs, beliefs + count); + std::vector result(count); + backend->CopyToHost(result.data(), beliefs, count); + return result; } void SimplePCNetwork::SetMuCacheThreshold(float threshold) noexcept @@ -171,7 +232,7 @@ namespace Deep void SimplePCNetwork::Compile() { #pragma omp parallel - { // Broadcast FTZ/DAZ hardware flags to ALL OpenMP worker threads + { _MM_SET_FLUSH_ZERO_MODE(_MM_FLUSH_ZERO_ON); _MM_SET_DENORMALS_ZERO_MODE(_MM_DENORMALS_ZERO_ON); } @@ -179,8 +240,21 @@ namespace Deep for (auto &layer : layers) total_floats_needed += layer->GetRequiredFloats(); - arena = std::make_unique(total_floats_needed, true); - for (auto &layer : layers) - layer->BindMemory(*arena); + backend->PrepareForBatchSize(batchSize); + + if (device == DeviceType::DEVICE_CPU) + { + cpuArena = std::make_unique(total_floats_needed, true); + for (auto &layer : layers) + layer->BindMemory(*cpuArena); + } +#if defined(DEEPITY_USE_CUDA) + else + { + gpuArena = std::make_unique(backend.get(), total_floats_needed); + for (auto &layer : layers) + layer->BindMemory(*gpuArena); + } +#endif } -} +} \ No newline at end of file diff --git a/src/backend/Backend.cpp b/src/backend/Backend.cpp new file mode 100644 index 0000000..ca7dc40 --- /dev/null +++ b/src/backend/Backend.cpp @@ -0,0 +1,28 @@ +#include +#include +#include + +#ifdef DEEPITY_USE_CUDA +#include +#endif + +namespace Deep +{ + std::unique_ptr CreateBackend(DeviceType device) + { + if (device == DeviceType::DEVICE_CPU) + { + return std::make_unique(); + } + else if (device == DeviceType::DEVICE_GPU) + { +#ifdef DEEPITY_USE_CUDA + return std::make_unique(); +#else + throw std::runtime_error("Deepity was compiled without CUDA support (-DDEEPITY_USE_CUDA=OFF). Cannot create CUDABackend."); +#endif + } + + throw std::invalid_argument("Unknown DeviceType requested from CreateBackend."); + } +} \ No newline at end of file diff --git a/src/backend/CPUBackend.cpp b/src/backend/CPUBackend.cpp new file mode 100644 index 0000000..d21a6aa --- /dev/null +++ b/src/backend/CPUBackend.cpp @@ -0,0 +1,221 @@ +#include +#include +#include +#include +#include + +#ifdef DEEPITY_USE_MKL +#include +#else +#include +#endif +#include + +namespace Deep +{ + float *CPUBackend::Allocate(size_t numFloats) + { +#ifdef _WIN32 + return (float *)_aligned_alloc(numFloats * sizeof(float), 16); +#else + return (float *)aligned_alloc(16, numFloats * sizeof(float)); +#endif + } + + void CPUBackend::Free(float *ptr) noexcept + { +#ifdef _WIN32 + _aligned_free(ptr); +#else + free(ptr); +#endif + } + + void CPUBackend::Zero(float *ptr, size_t numFloats) noexcept + { + memset(ptr, 0, numFloats * sizeof(float)); + } + + void CPUBackend::Copy(float *dst, const float *src, size_t numFloats) noexcept + { + memcpy(dst, src, numFloats * sizeof(float)); + } + + void CPUBackend::CopyFromHost(float *deviceDst, const float *hostSrc, size_t numFloats) noexcept + { + memcpy(deviceDst, hostSrc, numFloats * sizeof(float)); + } + + void CPUBackend::CopyToHost(float *hostDst, const float *deviceSrc, size_t numFloats) noexcept + { + memcpy(hostDst, deviceSrc, numFloats * sizeof(float)); + } + + void CPUBackend::RandomizeNormal(float *buf, size_t n, float mean, float stddev, uint32_t seed) noexcept + { + std::mt19937 seedGenerator(seed); + std::uniform_int_distribution seedDist; + + std::vector seeds(omp_get_max_threads()); + for (auto &s : seeds) + s = seedDist(seedGenerator); + +#pragma omp parallel if (!omp_in_parallel()) + { + std::mt19937 rng(seeds[omp_get_thread_num()]); + std::normal_distribution dist(mean, stddev); + +#pragma omp for + for (ptrdiff_t i = 0; i < (ptrdiff_t)n; ++i) + buf[i] = dist(rng); + } + } + + void CPUBackend::RandomizeUniform(float *buf, size_t n, float min, float max, uint32_t seed) noexcept + { + std::mt19937 seedGenerator(seed); + std::uniform_int_distribution seedDist; + + std::vector seeds(omp_get_max_threads()); + for (auto &s : seeds) + s = seedDist(seedGenerator); + +#pragma omp parallel if (!omp_in_parallel()) + { + std::mt19937 rng(seeds[omp_get_thread_num()]); + std::uniform_real_distribution dist(min, max); + +#pragma omp for + for (ptrdiff_t i = 0; i < (ptrdiff_t)n; ++i) + buf[i] = dist(rng); + } + } + + void CPUBackend::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 + { // Multithreading is the caller's job. + cblas_sgemm(CblasRowMajor, transA ? CblasTrans : CblasNoTrans, transB ? CblasTrans : CblasNoTrans, + M, N, K, alpha, A, lda, B, ldb, beta, C, ldc); + } + + void CPUBackend::SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept + { + memset(dst, 0, width * sizeof(float)); + for (size_t b = 0; b < batchSize; ++b) + cblas_saxpy(width, 1.0f, src + b * width, 1, dst, 1); + } + + void CPUBackend::Scale(float *buf, size_t n, float alpha) noexcept + { + cblas_sscal(n, alpha, buf, 1); + } + + void CPUBackend::AxpyInto(float *y, const float *x, size_t n, float alpha) noexcept + { + cblas_saxpy(n, alpha, x, 1, y, 1); + } + + void CPUBackend::AddBiasBroadcast(float *buf, const float *bias, size_t batchSize, size_t width) noexcept + { +#pragma omp parallel for schedule(static) if (batchSize > 4 && !omp_in_parallel()) + for (size_t b = 0; b < batchSize; ++b) + cblas_saxpy(width, 1.0f, bias, 1, buf + b * width, 1); + } + + void CPUBackend::Activation(ActivationType type, float *buf, size_t n) noexcept + { + To_Fn(type)(buf, n); + } + + void CPUBackend::ActivationInto(ActivationType type, float *dst, const float *src, size_t n) noexcept + { + cblas_scopy(n, src, 1, dst, 1); + To_Fn(type)(dst, n); + } + + void CPUBackend::ActivationDerivative(ActivationType type, float *buf, size_t n, bool activated) noexcept + { + To_dFn(type)(buf, n, activated); + } + + void CPUBackend::ActivationDerivativeInto(ActivationType type, float *dst, const float *src, size_t n) noexcept + { + To_dFn2(type)(dst, src, n); + } + + void CPUBackend::FusedStateUpdate(float *z, const float *feedback, const float *deriv, + const float *e, size_t n, float ir) noexcept + { // TODO: Unknown if this is how it should look? +#pragma omp parallel for schedule(static) if (n > 4 && !omp_in_parallel()) + for (size_t i = 0; i < n; ++i) + { + z[i] += ir * ((feedback[i] * deriv[i]) - e[i]); + } + } + + float CPUBackend::ComputeErrorAndEnergy(float *e, const float *z, const float *mu, size_t n) noexcept + { + float energy = 0.0f; + +#pragma omp parallel for reduction(+ : energy) schedule(static) if (n > 256 && !omp_in_parallel()) + for (size_t i = 0; i < n; ++i) + { + float err = z[i] - mu[i]; + e[i] = err; + energy += err * err; + } + + return 0.5f * energy; + } + + void CPUBackend::ComputeError(float *e, const float *z, const float *mu, size_t n) noexcept + { +#pragma omp parallel for schedule(static) if (n > 256 && !omp_in_parallel()) + for (size_t i = 0; i < n; ++i) + e[i] = z[i] - mu[i]; + } + + void CPUBackend::IncrementCounter(int *counter) noexcept { ++(*counter); } + + void CPUBackend::AdamStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, + float beta1, float beta2, float eps) noexcept + { + // Precompute bias correction + float beta1_t = 1.0f - Sleef_powf_u10(beta1, static_cast(*t)); + float beta2_t = 1.0f - Sleef_powf_u10(beta2, static_cast(*t)); + float step_size = *lr * Sleef_sqrtf(beta2_t) / beta1_t; + +#pragma omp parallel for schedule(static) if (n > 256 && !omp_in_parallel()) + for (size_t i = 0; i < n; ++i) + { + float g = grad[i]; + m[i] = beta1 * m[i] + (1.0f - beta1) * g; + v[i] = beta2 * v[i] + (1.0f - beta2) * (g * g); + + param[i] -= step_size * m[i] / (std::sqrt(v[i]) + eps); + } + } + + void CPUBackend::AdamWStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, float weightDecay, + float beta1, float beta2, float eps) noexcept + { + float beta1_t = 1.0f - Sleef_powf_u10(beta1, static_cast(*t)); + float beta2_t = 1.0f - Sleef_powf_u10(beta2, static_cast(*t)); + float step_size = *lr * Sleef_sqrtf(beta2_t) / beta1_t; + +#pragma omp parallel for schedule(static) if (n > 256 && !omp_in_parallel()) + for (size_t i = 0; i < n; ++i) + { + float g = grad[i]; + m[i] = beta1 * m[i] + (1.0f - beta1) * g; + v[i] = beta2 * v[i] + (1.0f - beta2) * (g * g); + + param[i] -= *lr * weightDecay * param[i]; + param[i] -= step_size * m[i] / (Sleef_sqrtf(v[i]) + eps); + } + } +} \ No newline at end of file diff --git a/src/backend/CUDABackend.cu b/src/backend/CUDABackend.cu new file mode 100644 index 0000000..0bffdbc --- /dev/null +++ b/src/backend/CUDABackend.cu @@ -0,0 +1,604 @@ +#include +#include +#include +#include + +#ifdef DEEPITY_USE_CUDA +#include + +namespace Deep +{ + CUDABackend::CUDABackend() + { + cudaStreamCreate(&this->stream); + cublasCreate(&this->handle); + cublasSetStream(this->handle, this->stream); + + // constexpr size_t WORKSPACE_SIZE = 4 * 1024 * 1024; + // cudaMalloc(&workspace, WORKSPACE_SIZE); + // cublasSetWorkspace(handle, workspace, WORKSPACE_SIZE); + } + + CUDABackend::~CUDABackend() + { + if (hasGraph) + { + cudaGraphExecDestroy(graphExec); + cudaGraphDestroy(graph); + } + if (workspace) + cudaFree(workspace); + cublasDestroy(this->handle); + cudaStreamDestroy(this->stream); + } + + void CUDABackend::BeginGraphCapture() noexcept + { + cudaError_t err = cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal); + if (err != cudaSuccess) + std::cerr << "cudaStreamBeginCapture failed: " << cudaGetErrorString(err) << "\n"; + } + + bool CUDABackend::EndGraphCapture() noexcept + { + cudaGraph_t newGraph; + cudaError_t err = cudaStreamEndCapture(stream, &newGraph); + if (err != cudaSuccess) + { + std::cerr << "cudaStreamEndCapture failed: " << cudaGetErrorString(err) << "\n"; + return false; + } + + if (hasGraph) + { + cudaGraphExecDestroy(graphExec); + cudaGraphDestroy(graph); + hasGraph = false; + } + + graph = newGraph; + err = cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0); + if (err != cudaSuccess) + { + std::cerr << "cudaGraphInstantiate failed: " << cudaGetErrorString(err) << "\n"; + cudaGraphDestroy(graph); + graph = nullptr; + return false; + } + + hasGraph = true; + return true; + } + + void CUDABackend::ReplayGraph() noexcept + { + if (!hasGraph) + { + std::cerr << "ReplayGraph() called before any graph was captured.\n"; + return; + } + cudaError_t err = cudaGraphLaunch(graphExec, stream); + if (err != cudaSuccess) + std::cerr << "cudaGraphLaunch failed: " << cudaGetErrorString(err) << "\n"; + } + + float *CUDABackend::Allocate(size_t numFloats) + { + float *ptr = nullptr; + if (cudaMalloc(&ptr, numFloats * sizeof(float)) != cudaSuccess) + { + std::cerr << "Could not allocate memory to CUDA backend.\n"; + return nullptr; + } + return ptr; + } + + void CUDABackend::Free(float *ptr) noexcept + { + if (cudaFree(ptr) != cudaSuccess) + std::cerr << "Could not free memory from CUDA backend.\n"; + } + + void CUDABackend::Zero(float *ptr, size_t numFloats) noexcept + { + cudaMemsetAsync(ptr, 0, numFloats * sizeof(float), stream); + } + + void CUDABackend::Copy(float *dst, const float *src, size_t numFloats) noexcept + { + cudaMemcpyAsync(dst, src, numFloats * sizeof(float), cudaMemcpyDefault, stream); + } + + void CUDABackend::CopyFromHost(float *deviceDst, const float *hostSrc, size_t numFloats) noexcept + { + cudaMemcpy(deviceDst, hostSrc, numFloats * sizeof(float), cudaMemcpyHostToDevice); + } + + void CUDABackend::CopyToHost(float *hostDst, const float *deviceSrc, size_t numFloats) noexcept + { + cudaMemcpy(hostDst, deviceSrc, numFloats * sizeof(float), cudaMemcpyDeviceToHost); + } + + __global__ void normal_generation(curandState *state, float *random_numbers, + size_t n, float mean, float stddev, uint32_t seed) + { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + curand_init(seed, static_cast(i), 0, &state[i]); + random_numbers[i] = mean + stddev * curand_normal(&state[i]); + } + } + + __global__ void uniform_generation(curandState *state, float *random_numbers, size_t n, float min, float range, uint32_t seed) + { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + curand_init(seed, static_cast(i), 0, &state[i]); + random_numbers[i] = min + curand_uniform(&state[i]) * range; + } + } + + void CUDABackend::RandomizeNormal(float *buf, size_t n, float mean, float stddev, uint32_t seed) noexcept + { + if (n == 0) + return; + + curandState *state = nullptr; + if (cudaMalloc(&state, n * sizeof(curandState)) != cudaSuccess) + return; + + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + normal_generation<<>>(state, buf, n, mean, stddev, seed); + cudaStreamSynchronize(stream); + cudaFree(state); + } + + void CUDABackend::RandomizeUniform(float *buf, size_t n, float min, float max, uint32_t seed) noexcept + { + if (n == 0) + return; + + curandState *state = nullptr; + cudaError_t mallocErr = cudaMalloc(&state, n * sizeof(curandState)); + if (mallocErr != cudaSuccess) + { + std::cerr << "curandState cudaMalloc failed for n=" << n + << " (" << n * sizeof(curandState) << " bytes): " + << cudaGetErrorString(mallocErr) << "\n"; + return; + } + + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + uniform_generation<<>>(state, buf, n, min, max - min, seed); + + cudaError_t launchErr = cudaGetLastError(); + if (launchErr != cudaSuccess) + std::cerr << "uniform_generation kernel launch failed: " << cudaGetErrorString(launchErr) << "\n"; + + cudaStreamSynchronize(stream); + cudaError_t syncErr = cudaGetLastError(); + if (syncErr != cudaSuccess) + std::cerr << "uniform_generation kernel execution failed: " << cudaGetErrorString(syncErr) << "\n"; + + cudaFree(state); + } + + __global__ void FillOnesKernel(float *buf, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + buf[i] = 1.0f; + } + + void CUDABackend::PrepareForBatchSize(size_t batchSize) noexcept + { + if (onesVector) + cudaFree(onesVector); + cudaMalloc(&onesVector, batchSize * sizeof(float)); + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((batchSize + BLOCK_SIZE - 1) / BLOCK_SIZE); + FillOnesKernel<<>>(onesVector, batchSize); + cudaStreamSynchronize(stream); + } + + 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 + { + cublasOperation_t cuTransA = transA ? CUBLAS_OP_T : CUBLAS_OP_N; + cublasOperation_t cuTransB = transB ? CUBLAS_OP_T : CUBLAS_OP_N; + cublasStatus_t status = cublasSgemm(handle, cuTransB, cuTransA, + N, M, K, &alpha, B, ldb, A, lda, &beta, C, ldc); + if (status != CUBLAS_STATUS_SUCCESS) + { + cudaError_t cudaErr = cudaGetLastError(); + std::cerr << "cublasSgemm status: " << status + << ", underlying cudaError: " << cudaErr + << " (" << cudaGetErrorString(cudaErr) << ")\n"; + } + } + + void CUDABackend::SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept + { + float alpha = 1.0f, beta = 0.0f; + cublasStatus_t status = cublasSgemv(handle, CUBLAS_OP_N, width, batchSize, + &alpha, src, width, onesVector, 1, &beta, dst, 1); + if (status != CUBLAS_STATUS_SUCCESS) + std::cerr << "cublasSgemv (SumRows) status: " << status << "\n"; + } + + void CUDABackend::Scale(float *buf, size_t n, float alpha) noexcept + { + if (cublasSscal(handle, n, &alpha, buf, 1) != CUBLAS_STATUS_SUCCESS) + std::cerr << "Failed to perform CUDA Scale.\n"; + } + + void CUDABackend::AxpyInto(float *y, const float *x, size_t n, float alpha) noexcept + { + if (cublasSaxpy(handle, n, &alpha, x, 1, y, 1)) + std::cerr << "Failed to perform CUDA Axpy.\n"; + } + + __global__ void AddBiasBroadcastKernel(float *buf, const float *bias, size_t batchSize, size_t width) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < batchSize * width) + buf[i] += bias[i % width]; + } + + void CUDABackend::AddBiasBroadcast(float *buf, const float *bias, size_t batchSize, size_t width) noexcept + { + constexpr int BLOCK_SIZE = 256; + size_t total = batchSize * width; + const int blocks = static_cast((total + BLOCK_SIZE - 1) / BLOCK_SIZE); + AddBiasBroadcastKernel<<>>(buf, bias, batchSize, width); + } + +#pragma region ACTIVATIONS_AND_KERNELS + + __global__ void ReluKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + dst[i] = fmaxf(0.0f, src[i]); + } + + __global__ void GeluKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float xi = src[i]; + float inner = MAGIC_GELU_1 * xi * (1.0f + MAGIC_GELU_2 * xi * xi); + + float t; +#if __CUDA_ARCH__ >= 800 + asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(inner)); +#else + t = tanhf(inner); +#endif + + dst[i] = 0.5f * xi * (1.0f + t); + } + } + + __global__ void tanhKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float y; + asm("tanh.approx.f32 %0, %1;" : "=f"(y) : "f"(src[i])); + dst[i] = y; + } + } + + __global__ void sigmoidKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float t; + asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(0.5f * src[i])); + dst[i] = fmaf(0.5f, t, 0.5f); + } + } + + __global__ void eSigmoidKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float s = src[i]; + dst[i] = 0.5f * (s / (1.0f + fabsf(s)) + 1.0f); + } + } + + __global__ void linearKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + dst[i] = src[i]; + } + + __global__ void dReluKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + dst[i] = (float)(src[i] > 0.0f); + } + + constexpr MAGIC_GELU_2_3 = 3.0f * MAGIC_GELU_2; + + __global__ void dGeluKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float x = src[i]; + float xsq = x * x; + float inner = MAGIC_GELU_1 * x * (1.0f + MAGIC_GELU_2 * xsq); + + float t; +#if __CUDA_ARCH__ >= 800 + asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(inner)); +#else + t = tanhf(inner); +#endif + + float gprime = MAGIC_GELU_1 * (1.0f + MAGIC_GELU_2_3 * xsq); + float term1 = 0.5f * (1.0f + t); + float term2 = 0.5f * x * gprime * (1.0f - t * t); + + dst[i] = term1 + term2; + } + } + + __global__ void dTanhKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float t; + asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(src[i])); + dst[i] = fmaf(-t, t, 1.0f); + } + } + + __global__ void dSigmoidKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float t; + asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(0.5f * src[i])); + dst[i] = 0.25f * fmaf(-t, t, 1.0f); + } + } + + __global__ void dSigmoidActivatedKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float s = src[i]; + dst[i] = fmaf(-s, s, s); + } + } + + __global__ void d_eSigmoidKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float a = 1.0f + fabsf(src[i]); + dst[i] = 0.5f / (a * a); + } + } + + __global__ void d_eSigmoidActivatedKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float s = src[i]; + dst[i] = 2.0f * fmaf(-s, s, s); + } + } + + __global__ void dLinearKernelInto(float *dst, const float *src, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + dst[i] = 1.0f; + } + + __global__ void FusedStateUpdateKernel(float *z, const float *feedback, const float *deriv, + const float *e, size_t n, float ir) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + z[i] += ir * ((feedback[i] * deriv[i]) - e[i]); + } + } + + __global__ void ComputeErrorKernel(float *e, const float *z, const float *mu, size_t n) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + e[i] = z[i] - mu[i]; + } + } + +#pragma endregion + + void CUDABackend::Activation(ActivationType type, float *buf, size_t n) noexcept + { + ActivationInto(type, buf, buf, n); + } + + void CUDABackend::ActivationInto(ActivationType type, float *dst, const float *src, size_t n) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + switch (type) + { + case ActivationType::RELU: + ReluKernelInto<<>>(dst, src, n); + break; + case ActivationType::GELU: + GeluKernelInto<<>>(dst, src, n); + case ActivationType::SIGMOID: + sigmoidKernelInto<<>>(dst, src, n); + break; + case ActivationType::eSIGMOID: + eSigmoidKernelInto<<>>(dst, src, n); + break; + case ActivationType::TANH: + tanhKernelInto<<>>(dst, src, n); + break; + case ActivationType::LINEAR: + linearKernelInto<<>>(dst, src, n); + break; + case ActivationType::NONE: + default: + break; + } + } + + void CUDABackend::ActivationDerivative(ActivationType type, float *buf, size_t n, bool activated) noexcept + { + ActivationDerivativeInto(type, buf, buf, n); + } + + void CUDABackend::ActivationDerivativeInto(ActivationType type, float *dst, const float *src, size_t n) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + switch (type) + { + case ActivationType::dRELU: + dReluKernelInto<<>>(dst, src, n); + break; + case ActivationType::dGELU: + dGeluKernelInto<<>>(dst, src, n); + case ActivationType::dSIGMOID: + dSigmoidKernelInto<<>>(dst, src, n); + break; + case ActivationType::d_eSIGMOID: + d_eSigmoidKernelInto<<>>(dst, src, n); + break; + case ActivationType::dTANH: + dTanhKernelInto<<>>(dst, src, n); + break; + case ActivationType::dLINEAR: + dLinearKernelInto<<>>(dst, src, n); + break; + case ActivationType::NONE: + default: + break; + } + } + + void CUDABackend::FusedStateUpdate(float *z, const float *feedback, const float *deriv, + const float *e, size_t n, float ir) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + FusedStateUpdateKernel<<>>(z, feedback, deriv, e, n, ir); + } + + float CUDABackend::ComputeErrorAndEnergy(float *e, const float *z, const float *mu, size_t n) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + ComputeErrorKernel<<>>(e, z, mu, n); + + float sum_of_squares = 0.0f; + cublasSdot(handle, n, e, 1, e, 1, &sum_of_squares); + + return 0.5f * sum_of_squares; + } + + void CUDABackend::ComputeError(float *e, const float *z, const float *mu, size_t n) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + ComputeErrorKernel<<>>(e, z, mu, n); + } + + __global__ void IncrementCounterKernel(int *counter) { *counter += 1; } + + __global__ void AdamStepKernel(float *param, const float *grad, float *m, float *v, + size_t n, const int *t_ptr, const float *lr_ptr, + float beta1, float beta2, float eps) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float beta1_t = 1.0f - powf(beta1, (float)(*t_ptr)); + float beta2_t = 1.0f - powf(beta2, (float)(*t_ptr)); + float step_size = *lr_ptr * sqrtf(beta2_t) / beta1_t; + + float g = grad[i]; + m[i] = beta1 * m[i] + (1.0f - beta1) * g; + v[i] = beta2 * v[i] + (1.0f - beta2) * (g * g); + param[i] -= step_size * m[i] / (sqrtf(v[i]) + eps); + } + } + + __global__ void AdamWStepKernel(float *param, const float *grad, float *m, float *v, + size_t n, const int *t_ptr, const float *lr_ptr, float weightDecay, + float beta1, float beta2, float eps) + { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + float beta1_t = 1.0f - powf(beta1, (float)(*t_ptr)); + float beta2_t = 1.0f - powf(beta2, (float)(*t_ptr)); + float step_size = *lr_ptr * sqrtf(beta2_t) / beta1_t; + + float g = grad[i]; + m[i] = beta1 * m[i] + (1.0f - beta1) * g; + v[i] = beta2 * v[i] + (1.0f - beta2) * (g * g); + param[i] -= *lr_ptr * weightDecay * param[i]; + param[i] -= step_size * m[i] / (sqrtf(v[i]) + eps); + } + } + + void CUDABackend::IncrementCounter(int *counter) noexcept + { + IncrementCounterKernel<<<1, 1, 0, stream>>>(counter); + } + + void CUDABackend::AdamStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, + float beta1, float beta2, float eps) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + AdamStepKernel<<>>(param, grad, m, v, n, t, lr, beta1, beta2, eps); + } + + void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, float weightDecay, + float beta1, float beta2, float eps) noexcept + { + constexpr int BLOCK_SIZE = 256; + const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + AdamWStepKernel<<>>(param, grad, m, v, n, t, lr, weightDecay, beta1, beta2, eps); + } +} + +#endif \ No newline at end of file diff --git a/src/backend/Tensor.cpp b/src/backend/Tensor.cpp new file mode 100644 index 0000000..89f35e3 --- /dev/null +++ b/src/backend/Tensor.cpp @@ -0,0 +1 @@ +#include \ No newline at end of file diff --git a/temp.py b/temp.py new file mode 100644 index 0000000..9181d83 --- /dev/null +++ b/temp.py @@ -0,0 +1,152 @@ +import numpy as np +import os +import sys +from time import perf_counter +from pydeepity import SimplePCN + +def load_full_mnist(): + import gzip + import urllib.request + print("Fetching canonical MNIST dataset (idx-ubyte, matching ngc-learn exactly)...") + base_url = "https://storage.googleapis.com/cvdf-datasets/mnist/" + files: dict[str, str] = { + "x_train": "train-images-idx3-ubyte.gz", + "y_train": "train-labels-idx1-ubyte.gz", + "x_test": "t10k-images-idx3-ubyte.gz", + "y_test": "t10k-labels-idx1-ubyte.gz" + } + data_dir = "./data" + os.makedirs(data_dir, exist_ok=True) + paths: dict[str, str] = {} + for key, fname in files.items(): + filepath = os.path.join(data_dir, fname) + paths[key] = filepath + if not os.path.exists(filepath): + print(f"Downloading {fname}...") + urllib.request.urlretrieve(base_url + fname, filepath) + with gzip.open(paths["x_train"], 'rb') as f: + X_train_raw = np.frombuffer(f.read(), np.uint8, offset=16).reshape(-1, 784) + with gzip.open(paths["x_test"], 'rb') as f: + X_test_raw = np.frombuffer(f.read(), np.uint8, offset=16).reshape(-1, 784) + with gzip.open(paths["y_train"], 'rb') as f: + y_train_labels = np.frombuffer(f.read(), np.uint8, offset=8) + with gzip.open(paths["y_test"], 'rb') as f: + y_test_labels = np.frombuffer(f.read(), np.uint8, offset=8) + X_train = X_train_raw.astype(np.float32) / 255.0 + X_test = X_test_raw.astype(np.float32) / 255.0 + eps = 0.001 + Y_train = np.full((y_train_labels.shape[0], 10), eps, dtype=np.float32) + Y_train[np.arange(y_train_labels.shape[0]), y_train_labels] = 1.0 - eps + return X_train, Y_train, X_test, y_test_labels + +def main() -> None: + SEED = int(sys.argv[1]) if len(sys.argv) > 1 else 7 + EPOCHS = int(sys.argv[2]) if len(sys.argv) > 2 else 15 + + X_train, Y_train, X_test, y_test_labels = load_full_mnist() + + BATCH_SIZE = 250 + STEPS = 30 + LR = 0.00373 + DECAY_RATE = 0.94 + LMBDA = 0.0 # Keep at 0.0! Weight decay breaks the PCN's symmetric feedback + + print(f"\nBuilding network (784->512->512->10), seed={SEED}...") + net = SimplePCN(batch_size=BATCH_SIZE, device="gpu") + + net.add_layer(784, 512, lr=LR, ir=0.091, act="linear", lmbda=LMBDA) + net.add_layer(512, 512, lr=LR, ir=0.091, act="sigmoid", lmbda=LMBDA) + net.add_layer(512, 10, lr=LR, ir=0.091, act="sigmoid", lmbda=LMBDA) + net.add_layer(10, 0, lr=LR, ir=0.091, act="linear", lmbda=LMBDA) + + # 1. Engage decoupled AdamW + net.set_optimizer("ADAMW") + net.compile() + + # 2. Engage C++ zero-energy bypass and mu caching + # net.set_mu_cache_threshold(0) + + net.randomize_weights("uniform(-0.3,0.3)") + + print(f"\nTraining with FORWARD-PROJECTION init: {EPOCHS} epochs, {STEPS} steps, " + f"lr={LR}, decay_rate={DECAY_RATE}...\n") + print("Reference (ngc-learn, real run): 26.91, 42.96, 60.12, 75.20, 84.68, 89.52,") + print(" 91.90, 93.45, 94.30, 94.80, 95.13, 95.38, 95.63, 95.74, 95.95 -- test 95.09%\n") + + rng = np.random.default_rng(SEED) + n_batches = len(X_train) // BATCH_SIZE + start_time = perf_counter() + epoch_accs = [] + + for epoch in range(EPOCHS): + current_lr = LR * (DECAY_RATE ** epoch) + net.set_learning_rate(current_lr) + + # Standard randomized batches + indices = rng.permutation(len(X_train)) + X_shuf, Y_shuf = X_train[indices], Y_train[indices] + + correct = 0 + total = 0 + epoch_energy = 0.0 + + for b in range(n_batches): + X_batch = X_shuf[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] + Y_batch = Y_shuf[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] + + # 3. Call C++ Native Loop (avoid Nanobind overhead) + compute_energy = (b == n_batches - 1) + energy = net.train_step_with_projection(X_batch, Y_batch, STEPS, compute_energy) + epoch_energy += energy + + + N_ACC_BATCHES = 10 + for b in range(min(N_ACC_BATCHES, n_batches)): + X_batch = X_shuf[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] + Y_batch = Y_shuf[b * BATCH_SIZE:(b + 1) * BATCH_SIZE] + + net.reset_state() + net.clamp_input(X_batch) + flat_beliefs = net.predict_with_projection(X_batch, STEPS) + terminal_beliefs = flat_beliefs.reshape(BATCH_SIZE, 10) + + pred = np.argmax(terminal_beliefs, axis=1) + true = np.argmax(Y_batch, axis=1) + correct += np.sum(pred == true) + total += BATCH_SIZE + + epoch_acc = 100.0 * correct / total + epoch_accs.append(epoch_acc) + avg_energy = epoch_energy / n_batches + elapsed = perf_counter() - start_time + print(f"Epoch {epoch+1}/{EPOCHS} | Time: {elapsed:.1f}s | Acc: {epoch_acc:.2f}% | Avg energy: {avg_energy:.4f}") + + train_time = perf_counter() - start_time + print(f"\nTraining complete in {train_time:.1f}s.") + + print("\nRunning final test evaluation (with forward-projection init)...") + correct = 0 + total = 0 + for i in range(0, len(X_test), BATCH_SIZE): + X_batch = X_test[i:i + BATCH_SIZE] + y_labels_batch = y_test_labels[i:i + BATCH_SIZE] + if len(X_batch) != BATCH_SIZE: + continue + + net.reset_state() + net.clamp_input(X_batch) + flat_beliefs = net.predict_with_projection(X_batch, STEPS) + terminal_beliefs = flat_beliefs.reshape(BATCH_SIZE, 10) + + pred_classes = np.argmax(terminal_beliefs, axis=1) + correct += np.sum(pred_classes == y_labels_batch) + total += BATCH_SIZE + + test_acc = 100.0 * correct / total + print(f"\n=== Result ===") + print(f"Deepity Peak Test Accuracy: {test_acc:.2f}% (ngc-learn: 95.09%)") + print(f"Train time: {train_time:.1f}s") + print(f"\nDeepity per-epoch: {[round(a,2) for a in epoch_accs]}") + +if __name__ == "__main__": + main() diff --git a/tests/tAddBiasBroadcastVerify.cpp b/tests/tAddBiasBroadcastVerify.cpp new file mode 100644 index 0000000..4cac8e5 --- /dev/null +++ b/tests/tAddBiasBroadcastVerify.cpp @@ -0,0 +1,83 @@ +/** + * @file tAddBiasBroadcastVerify.cpp + * @brief Verifies IComputeBackend::AddBiasBroadcast against an + * independently, hand-computed expected result -- written specifically + * because a live MNIST run regressed from ~97% accuracy to ~10% (exactly + * chance) immediately after this function replaced a per-batch-row loop + * of individual AxpyInto calls. Something in the new implementation is + * mathematically wrong; this isolates exactly what. + */ +#include +#include +#include +#include + +using namespace Deep; + +namespace +{ + bool CheckClose(const std::vector &actual, const std::vector &expected, + const char *backendName) + { + bool ok = true; + for (size_t i = 0; i < expected.size(); ++i) + { + if (std::fabs(actual[i] - expected[i]) > 1e-4f) + { + printf(" [%s] MISMATCH at index %zu: got %.4f, expected %.4f\n", + backendName, i, actual[i], expected[i]); + ok = false; + } + } + return ok; + } + + bool RunTest(DeviceType device, const char *backendName) + { + auto backend = CreateBackend(device); + + const size_t batchSize = 3, width = 4; + std::vector hBuf = {1, 2, 3, 4, 10, 20, 30, 40, 100, 200, 300, 400}; + std::vector hBias = {0.1f, 0.2f, 0.3f, 0.4f}; + std::vector expected = {1.1f, 2.2f, 3.3f, 4.4f, + 10.1f, 20.2f, 30.3f, 40.4f, + 100.1f, 200.2f, 300.3f, 400.4f}; + + Tensor buf(backend.get(), device, hBuf); + Tensor bias(backend.get(), device, hBias); + + backend->AddBiasBroadcast(buf.Data(), bias.Data(), batchSize, width); + + std::vector out; + buf.CopyToHost(out); + + bool ok = CheckClose(out, expected, backendName); + printf(" [%s] AddBiasBroadcast: %s (got [%.1f, %.1f, %.1f, %.1f, ...])\n", + backendName, ok ? "PASSED" : "FAILED", out[0], out[1], out[2], out[3]); + return ok; + } +} + +int main() +{ + bool allPassed = true; + + printf("--- CPUBackend ---\n"); + allPassed &= RunTest(DeviceType::DEVICE_CPU, "CPU"); + +#ifdef DEEPITY_USE_CUDA + printf("\n--- CUDABackend ---\n"); + allPassed &= RunTest(DeviceType::DEVICE_GPU, "CUDA"); +#else + printf("\n--- CUDABackend skipped (DEEPITY_USE_CUDA not defined) ---\n"); +#endif + + if (!allPassed) + { + printf("\nFAILED: AddBiasBroadcast produced incorrect results on at least one backend.\n"); + return 1; + } + + printf("\nPASSED: AddBiasBroadcast correct on all tested backends.\n"); + return 0; +} \ No newline at end of file diff --git a/tests/tCUDAFunctionsVerify.cpp b/tests/tCUDAFunctionsVerify.cpp new file mode 100644 index 0000000..d10b0fa --- /dev/null +++ b/tests/tCUDAFunctionsVerify.cpp @@ -0,0 +1,237 @@ +/** + * @file tCUDAFunctionsVerify.cpp + * @brief Verifies every IComputeBackend method not already covered by + * tMatMulVerify.cpp -- activations, derivatives, Scale, AxpyInto, + * FusedStateUpdate, ComputeErrorAndEnergy, AdamStep, AdamWStep -- on + * both CPUBackend and CUDABackend, against expected values computed + * independently (via a separate Python script, not derived from this + * codebase's own formulas). + * + * IMPORTANT PRECISION NOTE: CUDABackend's tanh/sigmoid kernels use the + * hardware-approximate `tanh.approx.f32` PTX instruction, trading + * precision for speed -- CPUBackend uses SLEEF's high-precision + * (u10 = <=1.0 ULP error) implementation. These will NOT match to tight + * precision even when both are correct. Functions using this instruction + * (SIGMOID, TANH, and their derivatives) use a loose tolerance (1e-3); + * everything else (exact arithmetic: RELU, LINEAR, eSIGMOID, Scale, + * AxpyInto, FusedStateUpdate, ComputeErrorAndEnergy, Adam/AdamW) uses a + * tight one (1e-4), since a loose match there would hide a real bug. + */ +#include +#include +#include +#include +#include + +using namespace Deep; + +namespace +{ + int g_failures = 0; + + void Check(const std::vector &actual, const std::vector &expected, + const char *testName, const char *backendName, float tolerance) + { + bool ok = true; + for (size_t i = 0; i < expected.size(); ++i) + { + if (std::fabs(actual[i] - expected[i]) > tolerance) + { + printf(" [%s / %s] MISMATCH at index %zu: got %.6f, expected %.6f (tol %.1e)\n", + backendName, testName, i, actual[i], expected[i], tolerance); + ok = false; + } + } + if (ok) + printf(" [%s / %s] PASSED\n", backendName, testName); + else + g_failures++; + } + + void CheckScalar(float actual, float expected, const char *testName, + const char *backendName, float tolerance) + { + Check({actual}, {expected}, testName, backendName, tolerance); + } + + void RunAllTests(DeviceType device, const char *backendName) + { + auto backend = CreateBackend(device); + const std::vector xs = {-2.0f, -1.0f, -0.5f, 0.0f, 0.5f, 1.0f, 2.0f}; + const size_t n = xs.size(); + + // --- Activations (independently computed via Python) --------- + { + Tensor t(backend.get(), device, xs); + backend->Activation(ActivationType::RELU, t.Data(), n); + std::vector out; + t.CopyToHost(out); + Check(out, {0, 0, 0, 0, 0.5f, 1.0f, 2.0f}, "relu", backendName, 1e-4f); + } + { + Tensor t(backend.get(), device, xs); + backend->Activation(ActivationType::SIGMOID, t.Data(), n); + std::vector out; + t.CopyToHost(out); + Check(out, {0.11920292f, 0.26894142f, 0.37754067f, 0.5f, 0.62245933f, 0.73105858f, 0.88079708f}, + "sigmoid", backendName, 1e-3f); // approx tanh-based on GPU + } + { + Tensor t(backend.get(), device, xs); + backend->Activation(ActivationType::eSIGMOID, t.Data(), n); + std::vector out; + t.CopyToHost(out); + Check(out, {0.16666667f, 0.25f, 0.33333333f, 0.5f, 0.66666667f, 0.75f, 0.83333333f}, + "e_sigmoid", backendName, 1e-4f); // exact arithmetic, no approx instruction + } + { + Tensor t(backend.get(), device, xs); + backend->Activation(ActivationType::TANH, t.Data(), n); + std::vector out; + t.CopyToHost(out); + Check(out, {-0.96402758f, -0.76159416f, -0.46211716f, 0.0f, 0.46211716f, 0.76159416f, 0.96402758f}, + "tanh", backendName, 1e-3f); // approx instruction on GPU + } + { + Tensor t(backend.get(), device, xs); + backend->Activation(ActivationType::LINEAR, t.Data(), n); + std::vector out; + t.CopyToHost(out); + Check(out, {-2, -1, -0.5f, 0, 0.5f, 1, 2}, "linear (identity)", backendName, 1e-4f); + } + + // --- Derivatives, raw-input (...Into) variants ---------------- + { + Tensor src(backend.get(), device, xs); + Tensor dst(backend.get(), device, n); + backend->ActivationDerivativeInto(ActivationType::dRELU, dst.Data(), src.Data(), n); + std::vector out; + dst.CopyToHost(out); + Check(out, {0, 0, 0, 0, 1.0f, 1.0f, 1.0f}, "dRelu", backendName, 1e-4f); + } + { + Tensor src(backend.get(), device, xs); + Tensor dst(backend.get(), device, n); + backend->ActivationDerivativeInto(ActivationType::dSIGMOID, dst.Data(), src.Data(), n); + std::vector out; + dst.CopyToHost(out); + Check(out, {0.10499359f, 0.19661193f, 0.23500371f, 0.25f, 0.23500371f, 0.19661193f, 0.10499359f}, + "dSigmoid", backendName, 1e-3f); + } + { + Tensor src(backend.get(), device, xs); + Tensor dst(backend.get(), device, n); + backend->ActivationDerivativeInto(ActivationType::d_eSIGMOID, dst.Data(), src.Data(), n); + std::vector out; + dst.CopyToHost(out); + Check(out, {0.05555556f, 0.125f, 0.22222222f, 0.5f, 0.22222222f, 0.125f, 0.05555556f}, + "d_eSigmoid", backendName, 1e-4f); + } + { + Tensor src(backend.get(), device, xs); + Tensor dst(backend.get(), device, n); + backend->ActivationDerivativeInto(ActivationType::dTANH, dst.Data(), src.Data(), n); + std::vector out; + dst.CopyToHost(out); + Check(out, {0.07065082f, 0.41997434f, 0.78644773f, 1.0f, 0.78644773f, 0.41997434f, 0.07065082f}, + "dTanh", backendName, 1e-3f); + } + { + Tensor src(backend.get(), device, xs); + Tensor dst(backend.get(), device, n); + backend->ActivationDerivativeInto(ActivationType::dLINEAR, dst.Data(), src.Data(), n); + std::vector out; + dst.CopyToHost(out); + Check(out, {1, 1, 1, 1, 1, 1, 1}, "dLinear", backendName, 1e-4f); + } + + // --- Scale: [1,2,3,4] *= 2.0 ----------------------------------- + { + Tensor t(backend.get(), device, std::vector{1, 2, 3, 4}); + backend->Scale(t.Data(), 4, 2.0f); + std::vector out; + t.CopyToHost(out); + Check(out, {2, 4, 6, 8}, "Scale", backendName, 1e-4f); + } + + // --- AxpyInto: y=[1,1,1,1] += 2.0 * x=[1,2,3,4] ----------------- + { + Tensor y(backend.get(), device, std::vector{1, 1, 1, 1}); + Tensor x(backend.get(), device, std::vector{1, 2, 3, 4}); + backend->AxpyInto(y.Data(), x.Data(), 4, 2.0f); + std::vector out; + y.CopyToHost(out); + Check(out, {3, 5, 7, 9}, "AxpyInto", backendName, 1e-4f); + } + + // --- FusedStateUpdate: z += ir*(feedback*deriv - e) ------------- + { + Tensor z(backend.get(), device, std::vector{1.0f, 2.0f}); + Tensor feedback(backend.get(), device, std::vector{0.5f, 1.0f}); + Tensor deriv(backend.get(), device, std::vector{2.0f, 0.5f}); + Tensor e(backend.get(), device, std::vector{0.1f, 0.2f}); + backend->FusedStateUpdate(z.Data(), feedback.Data(), deriv.Data(), e.Data(), 2, 0.1f); + std::vector out; + z.CopyToHost(out); + Check(out, {1.09f, 2.03f}, "FusedStateUpdate", backendName, 1e-4f); + } + + // --- ComputeErrorAndEnergy: e=z-mu, returns 0.5*sum(e^2) -------- + { + Tensor z(backend.get(), device, std::vector{3.0f, 5.0f}); + Tensor mu(backend.get(), device, std::vector{1.0f, 2.0f}); + Tensor e(backend.get(), device, 2); + float energy = backend->ComputeErrorAndEnergy(e.Data(), z.Data(), mu.Data(), 2); + std::vector eOut; + e.CopyToHost(eOut); + Check(eOut, {2.0f, 3.0f}, "ComputeErrorAndEnergy (e)", backendName, 1e-4f); + CheckScalar(energy, 6.5f, "ComputeErrorAndEnergy (energy)", backendName, 1e-4f); + } + + // --- AdamStep: param=1.0, grad=0.5, m=v=0, t=1 ------------------ + { + Tensor param(backend.get(), device, std::vector{1.0f}); + Tensor grad(backend.get(), device, std::vector{0.5f}); + Tensor m(backend.get(), device, 1); + Tensor v(backend.get(), device, 1); + backend->AdamStep(param.Data(), grad.Data(), m.Data(), v.Data(), 1, 1, 0.1f); + std::vector out; + param.CopyToHost(out); + CheckScalar(out[0], 0.9000000632455132f, "AdamStep", backendName, 1e-4f); + } + + // --- AdamWStep: same, plus weightDecay=0.01 --------------------- + { + Tensor param(backend.get(), device, std::vector{1.0f}); + Tensor grad(backend.get(), device, std::vector{0.5f}); + Tensor m(backend.get(), device, 1); + Tensor v(backend.get(), device, 1); + backend->AdamWStep(param.Data(), grad.Data(), m.Data(), v.Data(), 1, 1, 0.1f, 0.01f); + std::vector out; + param.CopyToHost(out); + CheckScalar(out[0], 0.8990000632455132f, "AdamWStep", backendName, 1e-4f); + } + } +} + +int main() +{ + printf("--- CPUBackend ---\n"); + RunAllTests(DeviceType::DEVICE_CPU, "CPU"); + +#ifdef DEEPITY_USE_CUDA + printf("\n--- CUDABackend ---\n"); + RunAllTests(DeviceType::DEVICE_GPU, "CUDA"); +#else + printf("\n--- CUDABackend skipped (DEEPITY_USE_CUDA not defined) ---\n"); +#endif + + if (g_failures > 0) + { + printf("\nFAILED: %d check(s) mismatched.\n", g_failures); + return 1; + } + + printf("\nPASSED: all backend functions verified correct.\n"); + return 0; +} \ No newline at end of file diff --git a/tests/tMatMulLargeAsymmetricVerify.cpp b/tests/tMatMulLargeAsymmetricVerify.cpp new file mode 100644 index 0000000..1f44b23 --- /dev/null +++ b/tests/tMatMulLargeAsymmetricVerify.cpp @@ -0,0 +1,107 @@ +/** + * @file tMatMulLargeAsymmetricVerify.cpp + * @brief Reproduces, in isolation, the exact GEMM shape that crashed + * DirectFeedbackUpdate() on Tiny ImageNet: transA=true, transB=true, + * M=1024 (nextSize), N=12288 (size), K=250 (batchSize) -- K much + * smaller than M/N, a shape tMatMulVerify.cpp's tiny (M=N=K=2-3) cases + * never exercised, and one likely to select a different cuBLAS internal + * kernel (the crash was inside ampere_sgemm_128x64_tt specifically). + * + * Differential test: CPU's GEMM path is already trusted (verified + * correct in tMatMulVerify.cpp and used throughout tonight's CPU runs), + * so any GPU disagreement at this specific scale isolates a real, + * scale-dependent bug rather than a general transpose-logic error. + */ +#include +#include +#include +#include +#include + +using namespace Deep; + +int main() +{ + // Exact real dimensions from DirectFeedbackUpdate()'s second GEMM + // on the Tiny ImageNet run (layer 0: size=12288, nextSize=1024, + // batchSize=250). + const int M = 1024; // nextSize + const int N = 12288; // size + const int K = 250; // batchSize + + std::mt19937 rng(42); + std::normal_distribution dist(0.0f, 1.0f); + + // proj stored [K, M] (batchSize, nextSize) -- transA=true means + // op(A) = A^T = [M, K]. + std::vector hProj(K * M); + for (auto &v : hProj) + v = dist(rng); + + // zF stored [K, N] (batchSize, size) -- transB=true means + // op(B) = B^T = [N, K]... wait, matching the real call: zF is + // passed as B with transB=true, stored [K, N] per the same + // batchSize-major convention as every other buffer in this codebase. + std::vector hZF(K * N); + for (auto &v : hZF) + v = dist(rng); + + // W stored [M, N] (nextSize, size), started at zero (beta=1.0 in + // the real call accumulates onto existing weights; zero here isolates + // just this GEMM's own contribution for comparison). + std::vector hWZero(M * N, 0.0f); + + float alpha = 0.5f; // arbitrary, matches fl/batchSize's role + float beta = 1.0f; + + auto RunOn = [&](DeviceType device, const char *name) -> std::vector + { + auto backend = CreateBackend(device); + + Tensor proj(backend.get(), device, hProj); + Tensor zF(backend.get(), device, hZF); + Tensor W(backend.get(), device, hWZero); + + backend->MatMul( + /*transA=*/true, /*transB=*/true, + M, N, K, + alpha, proj.Data(), M, + zF.Data(), N, + beta, W.Data(), N); + + std::vector result(M * N); + W.CopyToHost(result.data()); + printf("[%s] done\n", name); + return result; + }; + + printf("Running CPU reference...\n"); + std::vector cpuResult = RunOn(DeviceType::DEVICE_CPU, "CPU"); + + printf("Running GPU (this is where the real crash happened)...\n"); + std::vector gpuResult = RunOn(DeviceType::DEVICE_GPU, "GPU"); + + printf("Comparing...\n"); + int mismatches = 0; + float maxDiff = 0.0f; + for (size_t i = 0; i < cpuResult.size(); ++i) + { + float diff = std::fabs(cpuResult[i] - gpuResult[i]); + maxDiff = std::max(maxDiff, diff); + if (diff > 1e-2f) + mismatches++; + } + + printf("Max diff: %f, mismatches (>1e-2): %d / %zu\n", maxDiff, mismatches, cpuResult.size()); + + if (mismatches > 0) + { + printf("FAILED: GPU MatMul disagrees with CPU at this exact real-world scale.\n"); + return 1; + } + + printf("PASSED (though note: if this test PASSES but the real training run still\n"); + printf("crashes, the bug is NOT in MatMul itself -- look at buffer sizing/lifetime\n"); + printf("in DirectKPPCLayer's actual DirectFeedbackUpdate() call site instead.\n"); + return 0; +} \ No newline at end of file