Merge branch 'development' into issue-515-remove-const-parameters-primitive

d-kamath · web-flow · commit db9943866603 · 2023-10-14T08:52:06.000-07:00
diff --git a/Simulator/Core/EdgeIndexMap.h b/Simulator/Core/EdgeIndexMap.h
@@ -11,7 +11,7 @@
  *  Edge indices are stored in the edge forward map (forwardIndex), and
  *  the pointer and length of the vertex i's outgoing edge indices are specified
  *  by outgoingSynapse_begin[i] and edgeCount[i] respectively.
- *  The incoming edges list is used in calcSummationMapDevice() device function to
+ *  The incoming edges list is used in calcSummationPointDevice() device function to
  *  calculate sum of PSRs for each vertex simultaneously.
  *  The list also used in AllSpikingNeurons::advanceVertices() function to allow back propagation.
  *
diff --git a/Simulator/Core/GPUModel.cpp b/Simulator/Core/GPUModel.cpp
@@ -171,22 +171,22 @@ void GPUModel::advance()
 #endif   // PERFORMANCE_METRICS
 
    // calculate summation point
-   calcSummationMap();
+   calcSummationPoint();
 
 #ifdef PERFORMANCE_METRICS
    cudaLapTime(t_gpu_calcSummation);
 #endif   // PERFORMANCE_METRICS
 }
 
 /// Add psr of all incoming synapses to summation points.
-void GPUModel::calcSummationMap()
+void GPUModel::calcSummationPoint()
 {
    // CUDA parameters
    const int threadsPerBlock = 256;
    int blocksPerGrid
       = (Simulator::getInstance().getTotalVertices() + threadsPerBlock - 1) / threadsPerBlock;
 
-   calcSummationMapDevice<<<blocksPerGrid, threadsPerBlock>>>(
+   calcSummationPointDevice<<<blocksPerGrid, threadsPerBlock>>>(
       Simulator::getInstance().getTotalVertices(), allVerticesDevice_, synapseIndexMapDevice_,
       allEdgesDevice_);
 }
@@ -324,10 +324,10 @@ void GPUModel::copySynapseIndexMapHostToDevice(EdgeIndexMap &synapseIndexMapHost
 /// @param[in] synapseIndexMapDevice_  Pointer to forward map structures in device memory.
 /// @param[in] allEdgesDevice      Pointer to Synapse structures in device memory.
 __global__ void
-   calcSummationMapDevice(int totalVertices,
-                          AllSpikingNeuronsDeviceProperties *__restrict__ allVerticesDevice,
-                          const EdgeIndexMapDevice *__restrict__ synapseIndexMapDevice_,
-                          const AllSpikingSynapsesDeviceProperties *__restrict__ allEdgesDevice)
+   calcSummationPointDevice(int totalVertices,
+                            AllSpikingNeuronsDeviceProperties *__restrict__ allVerticesDevice,
+                            const EdgeIndexMapDevice *__restrict__ synapseIndexMapDevice_,
+                            const AllSpikingSynapsesDeviceProperties *__restrict__ allEdgesDevice)
 {
    // The usual thread ID calculation and guard against excess threads
    // (beyond the number of vertices, in this case).
@@ -355,7 +355,7 @@ __global__ void
          sum += allEdgesDevice->psr_[synIndex];
       }
       // Store summed PSR into this neuron's summation point
-      allVerticesDevice->summationMap_[idx] = sum;
+      allVerticesDevice->summationPoints_[idx] = sum;
    }
 }
 
diff --git a/Simulator/Core/GPUModel.h b/Simulator/Core/GPUModel.h
@@ -113,7 +113,7 @@ class GPUModel : public Model {
    virtual void deleteDeviceStruct(void **allVerticesDevice, void **allEdgesDevice);
 
    /// Add psr of all incoming synapses to summation points.
-   virtual void calcSummationMap();
+   virtual void calcSummationPoint();
 
    /// Pointer to device random noise array.
    float *randNoise_d;
@@ -161,8 +161,8 @@ void initMTGPU(unsigned int seed, unsigned int blocks, unsigned int threads, uns
 
 //! Calculate summation point.
 extern __global__ void
-   calcSummationMapDevice(int totalVertices,
-                          AllSpikingNeuronsDeviceProperties *__restrict__ allNeurnsDevice,
-                          const EdgeIndexMapDevice *__restrict__ synapseIndexMapDevice_,
-                          const AllSpikingSynapsesDeviceProperties *__restrict__ allEdgesDevice);
+   calcSummationPointDevice(int totalVertices,
+                            AllSpikingNeuronsDeviceProperties *__restrict__ allNeurnsDevice,
+                            const EdgeIndexMapDevice *__restrict__ synapseIndexMapDevice_,
+                            const AllSpikingSynapsesDeviceProperties *__restrict__ allEdgesDevice);
 #endif
diff --git a/Simulator/Edges/Neuro/AllSTDPSynapses.cpp b/Simulator/Edges/Neuro/AllSTDPSynapses.cpp
@@ -383,7 +383,7 @@ void AllSTDPSynapses::advanceEdge(BGSIZE iEdg, AllVertices &neurons)
    #ifdef USE_OMP
       #pragma omp atomic
    #endif
-   neurons.summationMap_[sumPointIndex] += psr;
+   neurons.summationPoints_[sumPointIndex] += psr;
    #ifdef USE_OMP
       //PAB: atomic above has implied flush (following statement generates error -- can't be member variable)
       //#pragma omp flush (summationPoint)
diff --git a/Simulator/Edges/Neuro/AllSpikingSynapses.cpp b/Simulator/Edges/Neuro/AllSpikingSynapses.cpp
@@ -286,7 +286,7 @@ void AllSpikingSynapses::advanceEdge(BGSIZE iEdg, AllVertices &neurons)
    #ifdef USE_OMP
       #pragma omp atomic #endif
    #endif
-   neurons.summationMap_[sumPointIndex] += psr;
+   neurons.summationPoints_[sumPointIndex] += psr;
    #ifdef USE_OMP
       //PAB: atomic above has implied flush (following statement generates error -- can't be member variable)
       //#pragma omp flush (summationPoint)
diff --git a/Simulator/Edges/Neuro/AllSynapsesDeviceFuncs.h b/Simulator/Edges/Neuro/AllSynapsesDeviceFuncs.h
@@ -44,7 +44,7 @@ extern __global__ void
 /// destination neurons.
 ///
 /// @param allEdgesDevice      Pointer to the Synapse structures in device memory.
-/// @param pSummationMap          Pointer to the summation point.
+/// @param pSummationPoint          Pointer to the summation point.
 /// @param deltaT                 The simulation time step size.
 /// @param weight                 Synapse weight.
 extern __global__ void initSynapsesDevice(int n, AllDSSynapsesDeviceProperties *allEdgesDevice,
diff --git a/Simulator/Edges/Neuro/AllSynapsesDeviceFuncs_d.cpp b/Simulator/Edges/Neuro/AllSynapsesDeviceFuncs_d.cpp
@@ -601,7 +601,7 @@ __global__ void updateSynapsesWeightsDevice(int numVertices, BGFLOAT deltaT, BGF
 /// destination neurons.
 ///
 /// @param allEdgesDevice      Pointer to the Synapse structures in device memory.
-/// @param pSummationMap          Pointer to the summation point.
+/// @param pSummationPoint          Pointer to the summation point.
 /// @param deltaT                 The simulation time step size.
 /// @param weight                 Synapse weight.
 __global__ void initSynapsesDevice(int n, AllDSSynapsesDeviceProperties *allEdgesDevice,
diff --git a/Simulator/Utils/Inputs/GpuSInputPoisson.cu b/Simulator/Utils/Inputs/GpuSInputPoisson.cu
@@ -82,7 +82,7 @@ void GpuSInputPoisson::inputStimulus(SimulationInfo* psi)
     advanceSpikingSynapsesDevice <<< blocksPerGrid, threadsPerBlock >>> ( edge_count, edgeIndexMapDevice, g_simulationStep, psi->deltaT, (AllSpikingSynapsesDeviceProperties*)allEdgesDevice );
 
     // update summation point
-    applyI2SummationMap <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, psi->pSummationMap, allEdgesDevice );
+    applyI2SummationMap <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, psi->pSummationPoint, allEdgesDevice );
 }
 
 /// Allocate GPU device memory and copy values
@@ -108,7 +108,7 @@ void GpuSInputPoisson::allocDeviceValues(IModel* model, SimulationInfo* psi, int
     const int threadsPerBlock = 256;
     int blocksPerGrid = ( vertex_count + threadsPerBlock - 1 ) / threadsPerBlock;
 
-    initSynapsesDevice <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, allEdgesDevice, psi->pSummationMap, psi->width, psi->deltaT, weight );
+    initSynapsesDevice <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, allEdgesDevice, psi->pSummationPoint, psi->width, psi->deltaT, weight );
 
     // allocate memory for curand global state
     HANDLE_ERROR( cudaMalloc ( &devStates_d, vertex_count * sizeof( curandState ) ) );
diff --git a/Simulator/Utils/Inputs/GpuSInputRegular.cu b/Simulator/Utils/Inputs/GpuSInputRegular.cu
@@ -71,7 +71,7 @@ void GpuSInputRegular::inputStimulus(SimulationInfo* psi)
 
     // add input to each summation point
     blocksPerGrid = ( vertex_count + threadsPerBlock - 1 ) / threadsPerBlock;
-    inputStimulusDevice <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, psi->pSummationMap, initValues_d, nShiftValues_d, nStepsInCycle, nStepsCycle, nStepsDuration );
+    inputStimulusDevice <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, psi->pSummationPoint, initValues_d, nShiftValues_d, nStepsInCycle, nStepsCycle, nStepsDuration );
     // update cycle count
     nStepsInCycle = (nStepsInCycle + 1) % nStepsCycle;
 }
diff --git a/Simulator/Utils/Inputs/HostSInputRegular.cpp b/Simulator/Utils/Inputs/HostSInputRegular.cpp
@@ -51,7 +51,7 @@ void HostSInputRegular::inputStimulus()
    for (int i = Simulator::getInstance().getTotalVertices() - 1; i >= 0; --i) {
       if ((nStepsInCycle >= nShiftValues[i])
           && (nStepsInCycle < (nShiftValues[i] + nStepsDuration) % nStepsCycle))
-         Simulator::getInstance().getLayout().getVertices().summationMap_[i] += values[i];
+         Simulator::getInstance().getLayout().getVertices().summationPoints_[i] += values[i];
    }
 
    // update cycle count
diff --git a/Simulator/Vertices/AllVertices.cpp b/Simulator/Vertices/AllVertices.cpp
@@ -32,12 +32,12 @@ void AllVertices::setupVertices()
 {
    size_ = Simulator::getInstance().getTotalVertices();
 #if defined(USE_GPU)
-   // We don't allocate memory for summationMap_ in CPU when building the GPU
+   // We don't allocate memory for summationPoints_ in CPU when building the GPU
    // implementation. This is to avoid misusing it in GPU code.
-   // summationMap_ = nullptr;
+   // summationPoints_ = nullptr;
 
 #else
-   summationMap_.assign(size_, 0);
+   summationPoints_.assign(size_, 0);
 
 #endif
 }
diff --git a/Simulator/Vertices/AllVertices.h b/Simulator/Vertices/AllVertices.h
@@ -76,7 +76,7 @@ class AllVertices {
    ///  apply (summed up) their PSRs (Post-Synaptic-Response).
    ///  On the next advance cycle, vertices add the values stored in their corresponding
    ///  summation points to their Vm and resets the summation points to zero
-   vector<BGFLOAT> summationMap_;
+   vector<BGFLOAT> summationPoints_;
 
 protected:
    ///  Total number of vertices.
@@ -144,6 +144,6 @@ struct AllVerticesDeviceProperties {
    ///  apply (summed up) their PSRs (Post-Synaptic-Response).
    ///  On the next advance cycle, vertices add the values stored in their corresponding
    ///  summation points to their Vm and resets the summation points to zero
-   BGFLOAT *summationMap_;
+   BGFLOAT *summationPoints_;
 };
 #endif   // defined(USE_GPU)
diff --git a/Simulator/Vertices/NG911/All911Vertices.cpp b/Simulator/Vertices/NG911/All911Vertices.cpp
@@ -279,5 +279,4 @@ void All911Vertices::advanceRESP(BGSIZE vertexIdx, All911Edges &edges911,
       LOG4CPLUS_DEBUG(vertexLogger_, "Responded to incident, type: " << incident->type);
    }
 }
-
 #endif
diff --git a/Simulator/Vertices/Neuro/AllIFNeurons_d.cpp b/Simulator/Vertices/Neuro/AllIFNeurons_d.cpp
@@ -49,7 +49,7 @@ void AllIFNeurons::allocDeviceStruct(AllIFNeuronsDeviceProperties &allVerticesDe
    HANDLE_ERROR(cudaMalloc((void **)&allVerticesDevice.hasFired_, count * sizeof(bool)));
    HANDLE_ERROR(
       cudaMalloc((void **)&allVerticesDevice.numStepsInRefractoryPeriod_, count * sizeof(int)));
-   HANDLE_ERROR(cudaMalloc((void **)&allVerticesDevice.summationMap_, count * sizeof(BGFLOAT)));
+   HANDLE_ERROR(cudaMalloc((void **)&allVerticesDevice.summationPoints_, count * sizeof(BGFLOAT)));
    HANDLE_ERROR(cudaMalloc((void **)&allVerticesDevice.spikeHistory_, count * sizeof(uint64_t *)));
 
    uint64_t *pSpikeHistory[count];
@@ -106,7 +106,7 @@ void AllIFNeurons::deleteDeviceStruct(AllIFNeuronsDeviceProperties &allVerticesD
    HANDLE_ERROR(cudaFree(allVerticesDevice.Vthresh_));
    HANDLE_ERROR(cudaFree(allVerticesDevice.hasFired_));
    HANDLE_ERROR(cudaFree(allVerticesDevice.numStepsInRefractoryPeriod_));
-   HANDLE_ERROR(cudaFree(allVerticesDevice.summationMap_));
+   HANDLE_ERROR(cudaFree(allVerticesDevice.summationPoints_));
    HANDLE_ERROR(cudaFree(allVerticesDevice.spikeHistory_));
 }
 
diff --git a/Simulator/Vertices/Neuro/AllIZHNeurons.cpp b/Simulator/Vertices/Neuro/AllIZHNeurons.cpp
@@ -211,7 +211,7 @@ void AllIZHNeurons::advanceNeuron(int index)
 {
    BGFLOAT &Vm = this->Vm_[index];
    BGFLOAT &Vthresh = this->Vthresh_[index];
-   BGFLOAT &summationPoint = this->summationMap_[index];
+   BGFLOAT &summationPoint = this->summationPoints_[index];
    BGFLOAT &I0 = this->I0_[index];
    BGFLOAT &Inoise = this->Inoise_[index];
    BGFLOAT &C1 = this->C1_[index];
diff --git a/Simulator/Vertices/Neuro/AllIZHNeurons_d.cpp b/Simulator/Vertices/Neuro/AllIZHNeurons_d.cpp
@@ -229,7 +229,7 @@ __global__ void advanceIZHNeuronsDevice(int totalVertices, int maxEdges, int max
       return;
 
    allVerticesDevice->hasFired_[idx] = false;
-   BGFLOAT &sp = allVerticesDevice->summationMap_[idx];
+   BGFLOAT &sp = allVerticesDevice->summationPoints_[idx];
    BGFLOAT &vm = allVerticesDevice->Vm_[idx];
    BGFLOAT &a = allVerticesDevice->Aconst_[idx];
    BGFLOAT &b = allVerticesDevice->Bconst_[idx];
diff --git a/Simulator/Vertices/Neuro/AllLIFNeurons.cpp b/Simulator/Vertices/Neuro/AllLIFNeurons.cpp
@@ -25,7 +25,7 @@ void AllLIFNeurons::advanceNeuron(int index)
 {
    BGFLOAT &Vm = this->Vm_[index];
    BGFLOAT &Vthresh = this->Vthresh_[index];
-   BGFLOAT &summationPoint = this->summationMap_[index];
+   BGFLOAT &summationPoint = this->summationPoints_[index];
    BGFLOAT &I0 = this->I0_[index];
    BGFLOAT &Inoise = this->Inoise_[index];
    BGFLOAT &C1 = this->C1_[index];
diff --git a/Simulator/Vertices/Neuro/AllLIFNeurons_d.cpp b/Simulator/Vertices/Neuro/AllLIFNeurons_d.cpp
@@ -95,7 +95,7 @@ __global__ void advanceLIFNeuronsDevice(int totalVertices, int maxEdges, int max
       return;
 
    allVerticesDevice->hasFired_[idx] = false;
-   BGFLOAT &sp = allVerticesDevice->summationMap_[idx];
+   BGFLOAT &sp = allVerticesDevice->summationPoints_[idx];
    BGFLOAT &vm = allVerticesDevice->Vm_[idx];
    BGFLOAT r_sp = sp;
    BGFLOAT r_vm = vm;
diff --git a/docs/Developer/ClassDiagrams/core.puml b/docs/Developer/ClassDiagrams/core.puml
@@ -28,7 +28,7 @@ class GPUModel {
 	+advance() : void
 	#allocDeviceStruct(void** allVerticesDevice, void** allEdgesDevice) : void
 	-allocSynapseImap(int count) : void
-	#calcSummationMap() : void
+	#calcSummationPoint() : void
 	+copyCPUtoGPU() : void
 	+copyGPUtoCPU() : void
 	+copySynapseIndexMapHostToDevice(EdgeIndexMap& synapseIndexMapHost, int numVertices) : void
@@ -76,8 +76,6 @@ class Simulator {
 	+getEpochDuration() : BGFLOAT {query}
 	+getMaxRate() : BGFLOAT {query}
 	-maxRate_ : BGFLOAT
-	+getPSummationMap() : BGFLOAT* {query}
-	-pSummationMap_ : BGFLOAT*
 	+{static} getInstance() : Simulator&
 	+getShort_timer() : Timer
 	+getTimer() : Timer
@@ -130,7 +128,6 @@ class Simulator {
 	+saveResults() : void {query}
 	+setConfigFileName(const string& fileName) : void
 	+setDeserializationFileName(const string& fileName) : void
-	+setPSummationMap(BGFLOAT* summationMap) : void
 	+setSerializationFileName(const string& fileName) : void
 	+setStimulusFileName(const string& fileName) : void
 	+setup() : void
diff --git a/docs/Developer/ClassDiagrams/diagrams/GraphittiClassDiagram.svg b/docs/Developer/ClassDiagrams/diagrams/GraphittiClassDiagram.svg
diff --git a/docs/Developer/ClassDiagrams/diagrams/GraphittiCoreClassDiagram.svg b/docs/Developer/ClassDiagrams/diagrams/GraphittiCoreClassDiagram.svg
diff --git a/docs/Developer/ClassDiagrams/diagrams/LayoutClassDiagram.svg b/docs/Developer/ClassDiagrams/diagrams/LayoutClassDiagram.svg
diff --git a/docs/Developer/ClassDiagrams/diagrams/VerticesClassDiagram.svg b/docs/Developer/ClassDiagrams/diagrams/VerticesClassDiagram.svg
diff --git a/docs/Developer/ClassDiagrams/graphitti.puml b/docs/Developer/ClassDiagrams/graphitti.puml
diff --git a/docs/Developer/ClassDiagrams/layout.puml b/docs/Developer/ClassDiagrams/layout.puml
diff --git a/docs/Developer/ClassDiagrams/vertices.puml b/docs/Developer/ClassDiagrams/vertices.puml
diff --git a/docs/RebuildNotes/GeneralNotes.md b/docs/RebuildNotes/GeneralNotes.md
diff --git a/docs/RebuildNotes/NeuronsNotes.md b/docs/RebuildNotes/NeuronsNotes.md

Original file line number	Diff line number	Diff line change
`@@ -11,7 +11,7 @@`
`11`	`11`	`* Edge indices are stored in the edge forward map (forwardIndex), and`
`12`	`12`	`* the pointer and length of the vertex i's outgoing edge indices are specified`
`13`	`13`	`* by outgoingSynapse_begin[i] and edgeCount[i] respectively.`
`14`		`- * The incoming edges list is used in calcSummationMapDevice() device function to`
	`14`	`+ * The incoming edges list is used in calcSummationPointDevice() device function to`
`15`	`15`	`* calculate sum of PSRs for each vertex simultaneously.`
`16`	`16`	`* The list also used in AllSpikingNeurons::advanceVertices() function to allow back propagation.`
`17`	`17`	`*`
Original file line number	Diff line number	Diff line change
`@@ -71,7 +71,7 @@ void GpuSInputRegular::inputStimulus(SimulationInfo* psi)`
`71`	`71`
`72`	`72`	`// add input to each summation point`
`73`	`73`	`blocksPerGrid = ( vertex_count + threadsPerBlock - 1 ) / threadsPerBlock;`
`74`		`- inputStimulusDevice <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, psi->pSummationMap, initValues_d, nShiftValues_d, nStepsInCycle, nStepsCycle, nStepsDuration );`
	`74`	`+ inputStimulusDevice <<< blocksPerGrid, threadsPerBlock >>> ( vertex_count, psi->pSummationPoint, initValues_d, nShiftValues_d, nStepsInCycle, nStepsCycle, nStepsDuration );`
`75`	`75`	`// update cycle count`
`76`	`76`	`nStepsInCycle = (nStepsInCycle + 1) % nStepsCycle;`
`77`	`77`	`}`
Original file line number	Diff line number	Diff line change
`@@ -51,7 +51,7 @@ void HostSInputRegular::inputStimulus()`
`51`	`51`	`for (int i = Simulator::getInstance().getTotalVertices() - 1; i >= 0; --i) {`
`52`	`52`	`if ((nStepsInCycle >= nShiftValues[i])`
`53`	`53`	`&& (nStepsInCycle < (nShiftValues[i] + nStepsDuration) % nStepsCycle))`
`54`		`- Simulator::getInstance().getLayout().getVertices().summationMap_[i] += values[i];`
	`54`	`+ Simulator::getInstance().getLayout().getVertices().summationPoints_[i] += values[i];`
`55`	`55`	`}`
`56`	`56`
`57`	`57`	`// update cycle count`