@@ -279,58 +279,6 @@ void GPUModel::allocEdgeIndexMap(int count)
279279 cudaMemcpyHostToDevice));
280280}
281281
282- // / Calculate the sum of synaptic input to each neuron.
283- // /
284- // / Calculate the sum of synaptic input to each neuron. One thread
285- // / corresponds to one neuron. Iterates sequentially through the
286- // / forward synapse index map (edgeIndexMapDevice_) to access only
287- // / existing synapses. Using this structure eliminates the need to skip
288- // / synapses that have undergone lazy deletion from the main
289- // / (allEdgesDevice) synapse structure. The forward map is
290- // / re-computed during each network restructure (once per epoch) to
291- // / ensure that all synapse pointers for a neuron are stored
292- // / contiguously.
293- // /
294- // / @param[in] totalVertices Number of vertices in the entire simulation.
295- // / @param[in,out] allVerticesDevice Pointer to Neuron structures in device memory.
296- // / @param[in] edgeIndexMapDevice_ Pointer to forward map structures in device memory.
297- // / @param[in] allEdgesDevice Pointer to Synapse structures in device memory.
298- __global__ void
299- calcSummationPointDevice (int totalVertices,
300- AllSpikingNeuronsDeviceProperties *__restrict__ allVerticesDevice,
301- const EdgeIndexMapDevice *__restrict__ edgeIndexMapDevice_,
302- const AllSpikingSynapsesDeviceProperties *__restrict__ allEdgesDevice)
303- {
304- // The usual thread ID calculation and guard against excess threads
305- // (beyond the number of vertices, in this case).
306- int idx = blockIdx.x * blockDim.x + threadIdx.x ;
307- if (idx >= totalVertices)
308- return ;
309-
310- // Number of incoming synapses
311- const BGSIZE synCount = edgeIndexMapDevice_->incomingEdgeCount_ [idx];
312- // Optimization: terminate thread if no incoming synapses
313- if (synCount != 0 ) {
314- // Index of start of this neuron's block of forward map entries
315- const int beginIndex = edgeIndexMapDevice_->incomingEdgeBegin_ [idx];
316- // Address of the start of this neuron's block of forward map entries
317- const BGSIZE *activeMapBegin = &(edgeIndexMapDevice_->incomingEdgeIndexMap_ [beginIndex]);
318- // Summed post-synaptic response (PSR)
319- BGFLOAT sum = 0.0 ;
320- // Index of the current incoming synapse
321- BGSIZE synIndex;
322- // Repeat for each incoming synapse
323- for (BGSIZE i = 0 ; i < synCount; i++) {
324- // Get index of current incoming synapse
325- synIndex = activeMapBegin[i];
326- // Fetch its PSR and add into sum
327- sum += allEdgesDevice->psr_ [synIndex];
328- }
329- // Store summed PSR into this neuron's summation point
330- allVerticesDevice->summationPoints_ [idx] = sum;
331- }
332- }
333-
334282// / Allocate and Copy CPU Synapse data to GPU.
335283void GPUModel::copyCPUtoGPU ()
336284{
0 commit comments