Change (#370): reduce complexity of core

Update core-structure form a square of synapse-blocks into a linear order ob blocks. This reduce the flexibility a bit, but also the complexity, which solves the problem of the exact definition of required memory and makes gpu-support much easier.
kitsudaiki · Jul 15, 2024 · 540a5c6 · 540a5c6
1 parent 863f37d
commit 540a5c6
Show file tree

Hide file tree

Showing 9 changed files with 89 additions and 112 deletions.
diff --git a/src/Hanami/src/core/cluster/objects.h b/src/Hanami/src/core/cluster/objects.h
@@ -46,9 +46,9 @@ class Cluster;
 #define UNINTI_POINT_32 0x0FFFFFFF
 
 // network-predefines
-#define SYNAPSES_PER_SYNAPSESECTION 64
-#define NUMBER_OF_SYNAPSESECTION 64
-#define NEURONS_PER_NEURONBLOCK 64
+#define SYNAPSES_PER_SYNAPSESECTION 128
+#define NUMBER_OF_SYNAPSESECTION 512
+#define NEURONS_PER_NEURONBLOCK 128
 #define POSSIBLE_NEXT_AXON_STEP 80
 #define NUMBER_OF_POSSIBLE_NEXT 86
 #define NUMBER_OF_OUTPUT_CONNECTIONS 7
@@ -165,7 +165,7 @@ struct SynapseSection {
 
     SynapseSection() { std::fill_n(synapses, SYNAPSES_PER_SYNAPSESECTION, Synapse()); }
 };
-static_assert(sizeof(SynapseSection) == 1024);
+static_assert(sizeof(SynapseSection) == 2048);
 
 //==================================================================================================
 
@@ -174,7 +174,7 @@ struct SynapseBlock {
 
     SynapseBlock() { std::fill_n(sections, NUMBER_OF_SYNAPSESECTION, SynapseSection()); }
 };
-static_assert(sizeof(SynapseBlock) == 64 * 1024);
+static_assert(sizeof(SynapseBlock) == 512 * 2048);
 
 //==================================================================================================
 
@@ -224,7 +224,7 @@ struct NeuronBlock {
 
     NeuronBlock() { std::fill_n(neurons, NEURONS_PER_NEURONBLOCK, Neuron()); }
 };
-static_assert(sizeof(NeuronBlock) == 2048);
+static_assert(sizeof(NeuronBlock) == 4096);
 
 //==================================================================================================
 
@@ -289,7 +289,7 @@ struct ConnectionBlock {
 
     ConnectionBlock() { std::fill_n(connections, NUMBER_OF_SYNAPSESECTION, SynapseConnection()); }
 };
-static_assert(sizeof(ConnectionBlock) == 1544);
+static_assert(sizeof(ConnectionBlock) == 12296);
 
 //==================================================================================================
 
@@ -315,9 +315,9 @@ struct HexagonHeader {
     uint32_t hexagonId = UNINIT_STATE_32;
     bool isInputHexagon = false;
     bool isOutputHexagon = false;
-    uint8_t padding1[2];
+    uint8_t padding[2];
+    uint32_t numberOfFreeSections = 0;
     uint32_t dimX = 0;
-    uint32_t dimY = 0;
     Hanami::Position hexagonPos;
 
     bool operator==(HexagonHeader& rhs)
@@ -334,9 +334,6 @@ struct HexagonHeader {
         if (dimX != rhs.dimX) {
             return false;
         }
-        if (dimY != rhs.dimY) {
-            return false;
-        }
         if (hexagonPos != rhs.hexagonPos) {
             return false;
         }

diff --git a/src/Hanami/src/core/io/checkpoint/io_interface.cpp b/src/Hanami/src/core/io/checkpoint/io_interface.cpp
@@ -531,7 +531,7 @@ IO_Interface::checkHexagonEntry(const HexagonEntry& hexagonEntry)
     // check size against dimentsions in hexagon-header
     const uint64_t numberOfConnectionBlocks
         = hexagonEntry.numberOfConnectionBytes / (sizeof(ConnectionBlock) + sizeof(SynapseBlock));
-    if (numberOfConnectionBlocks != hexagonEntry.header.dimX * hexagonEntry.header.dimY) {
+    if (numberOfConnectionBlocks != hexagonEntry.header.dimX) {
         return false;
     }
 

diff --git a/src/Hanami/src/core/io/checkpoint/io_interface.h b/src/Hanami/src/core/io/checkpoint/io_interface.h
@@ -34,7 +34,7 @@ namespace Hanami
 struct DataBuffer;
 }
 
-#define LOCAL_BUFFER_SIZE 128 * 1024
+#define LOCAL_BUFFER_SIZE 1024 * 4096
 
 class IO_Interface
 {
@@ -95,6 +95,7 @@ class IO_Interface
     template <typename T>
     inline bool addObjectToLocalBuffer(T* data, Hanami::ErrorContainer& error)
     {
+        assert(sizeof(T) < LOCAL_BUFFER_SIZE);
         if (sizeof(T) + m_localBuffer.size > LOCAL_BUFFER_SIZE) {
             if (writeFromLocalBuffer(m_localBuffer, error) == false) {
                 error.addMessage("Failed to write local buffer to target");

diff --git a/src/Hanami/src/core/processing/cluster_resize.h b/src/Hanami/src/core/processing/cluster_resize.h
@@ -82,38 +82,24 @@ inline void
 resizeConnections(Hexagon* targetHexagon)
 {
     const uint32_t dimXold = targetHexagon->header.dimX;
-    const uint32_t dimYold = targetHexagon->header.dimY;
-    int32_t x, y = 0;
 
     targetHexagon->header.dimX++;
-    targetHexagon->header.dimY++;
 
     // resize list
-    targetHexagon->connectionBlocks.resize(targetHexagon->header.dimX * targetHexagon->header.dimY);
+    targetHexagon->connectionBlocks.resize(targetHexagon->header.dimX);
 
     // if there was no scaling in x-dimension, then no re-ordering necessary
     if (targetHexagon->header.dimX == dimXold) {
         return;
     }
-
-    LOG_DEBUG("resized connection-Block: " + std::to_string(dimXold) + ":" + std::to_string(dimYold)
-              + " -> " + std::to_string(targetHexagon->header.dimX) + ":"
-              + std::to_string(targetHexagon->header.dimY));
-    uint32_t newPos = 0;
-    uint32_t oldPos = 0;
+    LOG_DEBUG("resized connection-Block: " + std::to_string(dimXold) + " -> "
+              + std::to_string(targetHexagon->header.dimX));
 
     // update content of list for the new size
-    for (y = dimYold - 1; y >= 1; --y) {
-        for (x = dimXold - 1; x >= 0; --x) {
-            newPos = (y * targetHexagon->header.dimX) + x;
-            oldPos = (y * dimXold) + x;
-
-            targetHexagon->connectionBlocks[newPos] = targetHexagon->connectionBlocks[oldPos];
-            targetHexagon->connectionBlocks[oldPos] = ConnectionBlock();
-        }
-    }
+    targetHexagon->connectionBlocks[targetHexagon->header.dimX - 1] = ConnectionBlock();
 
     targetHexagon->neuronBlocks.resize(targetHexagon->header.dimX);
+    targetHexagon->header.numberOfFreeSections += NUMBER_OF_SYNAPSESECTION;
 }
 
 /**
@@ -146,12 +132,20 @@ createNewSection(Cluster& cluster,
     Hexagon* targetHexagon = &cluster.hexagons[targetHexagonId];
 
     // get target-connection
+    if (targetHexagon->header.numberOfFreeSections < NUMBER_OF_SYNAPSESECTION / 2) {
+        resizeConnections(targetHexagon);
+    }
     SynapseConnection* targetConnection = searchTargetInHexagon(targetHexagon, synapseBlockBuffer);
     if (targetConnection == nullptr) {
-        resizeConnections(targetHexagon);
-        targetConnection = searchTargetInHexagon(targetHexagon, synapseBlockBuffer);
-        targetHexagon->wasResized = true;
+        Hanami::ErrorContainer error;
+        error.addMessage("no target-section found, even there should be sill "
+                         + std::to_string(targetHexagon->header.numberOfFreeSections)
+                         + " available");
+        LOG_ERROR(error);
+        return false;
     }
+    targetHexagon->header.numberOfFreeSections--;
+    targetHexagon->wasResized = true;
 
     // initialize new connection
     targetConnection->origin = sourceLocPtr;

diff --git a/src/Hanami/src/core/processing/cpu/backpropagation.h b/src/Hanami/src/core/processing/cpu/backpropagation.h
@@ -128,34 +128,29 @@ backpropagateConnections(Hexagon* hexagon,
     NeuronBlock* targetNeuronBlock = nullptr;
     ConnectionBlock* connectionBlock = nullptr;
     SynapseSection* synapseSection = nullptr;
-    const uint32_t dimY = hexagon->header.dimY;
     const uint32_t dimX = hexagon->header.dimX;
     SourceLocation sourceLoc;
-    uint32_t c = 0;
-    uint32_t i = 0;
 
     if (blockId >= dimX) {
         return;
     }
 
-    for (c = blockId * dimY; c < (blockId * dimY) + dimY; ++c) {
-        assert(c < hexagon->connectionBlocks.size());
-        connectionBlock = &hexagon->connectionBlocks[c];
+    assert(blockId < hexagon->connectionBlocks.size());
+    connectionBlock = &hexagon->connectionBlocks[blockId];
 
-        for (i = 0; i < NUMBER_OF_SYNAPSESECTION; ++i) {
-            connection = &connectionBlock->connections[i];
+    for (uint32_t i = 0; i < NUMBER_OF_SYNAPSESECTION; ++i) {
+        connection = &connectionBlock->connections[i];
 
-            if (connection->origin.blockId == UNINIT_STATE_16) {
-                continue;
-            }
+        if (connection->origin.blockId == UNINIT_STATE_16) {
+            continue;
+        }
 
-            synapseSection = &synapseBlocks[connectionBlock->targetSynapseBlockPos].sections[i];
-            sourceLoc = getSourceNeuron(connection->origin, hexagons);
+        synapseSection = &synapseBlocks[connectionBlock->targetSynapseBlockPos].sections[i];
+        sourceLoc = getSourceNeuron(connection->origin, hexagons);
 
-            targetNeuronBlock = &hexagon->neuronBlocks[(c / hexagon->header.dimY)];
+        targetNeuronBlock = &hexagon->neuronBlocks[blockId];
 
-            backpropagateSection(synapseSection, connection, targetNeuronBlock, sourceLoc.neuron);
-        }
+        backpropagateSection(synapseSection, connection, targetNeuronBlock, sourceLoc.neuron);
     }
 }
 

diff --git a/src/Hanami/src/core/processing/cpu/processing.h b/src/Hanami/src/core/processing/cpu/processing.h
@@ -183,49 +183,44 @@ processSynapses(Cluster& cluster, Hexagon* hexagon, const uint32_t blockId)
     SynapseSection* section = nullptr;
     Hexagon* sourceHexagon = nullptr;
     uint32_t randomeSeed = rand();
-    const uint32_t dimY = hexagon->header.dimY;
     const uint32_t dimX = hexagon->header.dimX;
     SourceLocation sourceLoc;
-    uint32_t c = 0;
-    uint32_t i = 0;
 
     bool inputConnected = false;
 
     if (blockId >= dimX) {
         return;
     }
 
-    for (c = blockId * dimY; c < (blockId * dimY) + dimY; ++c) {
-        assert(c < hexagon->connectionBlocks.size());
-        connectionBlock = &hexagon->connectionBlocks[c];
+    assert(blockId < hexagon->connectionBlocks.size());
+    connectionBlock = &hexagon->connectionBlocks[blockId];
 
-        for (i = 0; i < NUMBER_OF_SYNAPSESECTION; ++i) {
-            scon = &connectionBlock->connections[i];
-            if (scon->origin.blockId == UNINIT_STATE_16) {
-                continue;
-            }
+    for (uint32_t i = 0; i < NUMBER_OF_SYNAPSESECTION; ++i) {
+        scon = &connectionBlock->connections[i];
+        if (scon->origin.blockId == UNINIT_STATE_16) {
+            continue;
+        }
 
-            inputConnected = scon->origin.isInput;
-            sourceLoc = getSourceNeuron(scon->origin, &cluster.hexagons[0]);
+        inputConnected = scon->origin.isInput;
+        sourceLoc = getSourceNeuron(scon->origin, &cluster.hexagons[0]);
 
-            if (sourceLoc.neuron->active == 0) {
-                continue;
-            }
+        if (sourceLoc.neuron->active == 0) {
+            continue;
+        }
 
-            section = &synapseBlocks[connectionBlock->targetSynapseBlockPos].sections[i];
-            targetNeuronBlock = &neuronBlocks[(c / hexagon->header.dimY)];
-            randomeSeed += (c * NUMBER_OF_SYNAPSESECTION) + i;
+        section = &synapseBlocks[connectionBlock->targetSynapseBlockPos].sections[i];
+        targetNeuronBlock = &neuronBlocks[blockId];
+        randomeSeed += (blockId * NUMBER_OF_SYNAPSESECTION) + i;
 
-            synapseProcessingBackward<doTrain>(cluster,
-                                               section,
-                                               scon,
-                                               targetNeuronBlock,
-                                               sourceLoc.neuron,
-                                               scon->origin,
-                                               clusterSettings,
-                                               inputConnected,
-                                               randomeSeed);
-        }
+        synapseProcessingBackward<doTrain>(cluster,
+                                           section,
+                                           scon,
+                                           targetNeuronBlock,
+                                           sourceLoc.neuron,
+                                           scon->origin,
+                                           clusterSettings,
+                                           inputConnected,
+                                           randomeSeed);
     }
 }
 

diff --git a/src/Hanami/src/core/processing/cpu/reduction.h b/src/Hanami/src/core/processing/cpu/reduction.h
@@ -86,39 +86,34 @@ reduceConnections(Hexagon* hexagon,
     Hexagon* sourceHexagon = nullptr;
     ConnectionBlock* connectionBlock = nullptr;
     SynapseSection* synapseSection = nullptr;
-    const uint32_t dimY = hexagon->header.dimY;
     const uint32_t dimX = hexagon->header.dimX;
-    uint32_t c = 0;
-    uint32_t i = 0;
 
     if (blockId >= dimX) {
         return;
     }
 
-    for (c = blockId * dimY; c < (blockId * dimY) + dimY; ++c) {
-        connectionBlock = &hexagon->connectionBlocks[c];
+    connectionBlock = &hexagon->connectionBlocks[blockId];
 
-        for (i = 0; i < NUMBER_OF_SYNAPSESECTION; ++i) {
-            connection = &connectionBlock->connections[i];
-            if (connection->origin.blockId == UNINIT_STATE_16) {
-                continue;
-            }
+    for (uint32_t i = 0; i < NUMBER_OF_SYNAPSESECTION; ++i) {
+        connection = &connectionBlock->connections[i];
+        if (connection->origin.blockId == UNINIT_STATE_16) {
+            continue;
+        }
 
-            synapseSection = &synapseBlocks[connectionBlock->targetSynapseBlockPos].sections[i];
-            sourceHexagon = &hexagons[connection->origin.hexagonId];
-            sourceNeuronBlock = &sourceHexagon->neuronBlocks[connection->origin.blockId];
-            sourceNeuron = &sourceNeuronBlock->neurons[connection->origin.neuronId];
+        synapseSection = &synapseBlocks[connectionBlock->targetSynapseBlockPos].sections[i];
+        sourceHexagon = &hexagons[connection->origin.hexagonId];
+        sourceNeuronBlock = &sourceHexagon->neuronBlocks[connection->origin.blockId];
+        sourceNeuron = &sourceNeuronBlock->neurons[connection->origin.neuronId];
 
-            // if section is complete empty, then erase it
-            if (reduceSection(synapseSection) == false) {
-                // initialize the creation of a new section
-                sourceNeuron->isNew = 1;
-                sourceNeuron->newLowerBound = connection->lowerBound;
+        // if section is complete empty, then erase it
+        if (reduceSection(synapseSection) == false) {
+            // initialize the creation of a new section
+            sourceNeuron->isNew = 1;
+            sourceNeuron->newLowerBound = connection->lowerBound;
 
-                // mark current connection as available again
-                connection->origin.blockId = UNINIT_STATE_16;
-                connection->origin.neuronId = UNINIT_STATE_8;
-            }
+            // mark current connection as available again
+            connection->origin.blockId = UNINIT_STATE_16;
+            connection->origin.neuronId = UNINIT_STATE_8;
         }
     }
 }

diff --git a/src/Hanami/src/core/processing/cuda/backpropagation.cu b/src/Hanami/src/core/processing/cuda/backpropagation.cu
@@ -173,11 +173,11 @@ backpropagation_CUDA(CudaClusterPointer* gpuPointer,
         backpropagateNeurons<<<hexagon->header.dimX, 64>>>(
                 gpuPointer->hexagonPointer[hexagonId].neuronBlocks);
 
-        backpropagateConnections<<<hexagon->header.dimX * hexagon->header.dimY, 64>>>(
+        backpropagateConnections<<<hexagon->header.dimX, 64>>>(
                 gpuPointer->hexagonPointer[hexagonId].neuronBlocks,
                 gpuPointer->synapseBlocks,
                 gpuPointer->hexagonPointer[hexagonId].connectionBlocks,
-                hexagon->header.dimY);
+                42);
 
         // copy neurons back to host
         cudaMemcpy(&hexagon->neuronBlocks[0],