CARLsim is an efficient, easy-to-use, GPU-accelerated software framework for simulating large-scale spiking neural network (SNN) models with a high degree of biological detail.
Issue: during loops that cycle through neuron ids, some neuron ids are incorrectly processed.
Scope: possibly only affects GPU code.
Example:
File: snn_gpu_module.cu
Function: kernel_STPUpdateAndDecayConductances()
Detailed reporting of the error in the code:
In the loop:
for (int bufPos = blockIdx.x; bufPos < totBuffers; bufPos += gridDim.x)
the nid, lastId, and grpId are indentified.
Background info:
Every neuron buffer group has a nid, lastId, and grpId. This is explained in code comments “static thread load allocation”. The purpose of this is to assign groups of neurons that can be processed by each GPU block that has a specific number of threads. This is known as “static thread load”. GPU multithreading, during each processing step, will loop through all its blocks and threads. Assigning neurons to threads in buffer groups allows for looping though threads to assign work of a specific thread to a specific neuron.
Because of this code design, each neuron group is split up into subgroups known as buffer groups. For instance, if a neuron group has 500 neurons and 128 threads are available, the number of buffer groups will be:
500/128=3.91 which is rounded up to 4.
in code: (int) ceil(1.0f * groupConfigs[netId][lGrpId].numN / bufSize);
For that neuron group, assuming it is the 1st neuron group in the sim, it will have buffer groups:
firstNid | lastId | grpId | bufferId
0, 127, 0, 0
128, 255, 0, 1
255, 383, 0, 2
384, 499, 0, 3
In the earlier mentioned loop, nid (first neuron id in buffer group), lastId, and grpId are iterated through.
The error that is occurring is when IS_REGULAR_NEURON() attempts to identify if the nid input to the function is a correct neuron index, in the sense that it represents a real neuron index and not a false neuron index. The function IS_REGULAR_NEURON() does this by checking if the nid is less than the total number of “numNReg + numNPois” which are all regular and poisson neurons. However, what IS_REGULAR_NEURON() fails to take into account is the current buffer’s lastId.
Let’s say there are 2 neuron groups with 2 neurons in each of them. The first group would have neuron id 0 and 1 and the second group has neuron ids 2 and 3. For the first buffer position in the earlier mentioned loop, neuron ids 0, 1, 2 and 3 will be passed to the IS_REGULAR_NEURON() function. However, only neurons 0 and 1 are real neuron indices because in the first buffer the lastId is 2 meaning +2 from the firstNid of 0. IS_REGULAR_NEURON() finds that in the first buffer nids 0-3 are less than the total number of neurons in the sim and therefore reports them as real indices. What is missing here is that IS_REGULAR_NEURON() should evaluate if the nids are within firstNid to lastId of the current buffer. That would eliminate nids 2 and 3 from being reported as real neurons in the first buffer processed in the loop.
Subsequently, in the STP processing for example, nids 0 and 1 are processed twice erroneously when applying the “*= configs.dAMPA” operation. This decays the AMPA synapse current signal too fast. This same IS_REGULAR_NEURON() loop error possibly affects other operations in the GPU code as well. This error is possibly limited to GPU code because of how the threads are processed in buffers for this code. If the CPU code does similar processing it may also be an error there.
Issue: during loops that cycle through neuron ids, some neuron ids are incorrectly processed. Scope: possibly only affects GPU code. Example: File: snn_gpu_module.cu Function: kernel_STPUpdateAndDecayConductances()
Detailed reporting of the error in the code: In the loop: for (int bufPos = blockIdx.x; bufPos < totBuffers; bufPos += gridDim.x) the nid, lastId, and grpId are indentified.
Background info: Every neuron buffer group has a nid, lastId, and grpId. This is explained in code comments “static thread load allocation”. The purpose of this is to assign groups of neurons that can be processed by each GPU block that has a specific number of threads. This is known as “static thread load”. GPU multithreading, during each processing step, will loop through all its blocks and threads. Assigning neurons to threads in buffer groups allows for looping though threads to assign work of a specific thread to a specific neuron.
Because of this code design, each neuron group is split up into subgroups known as buffer groups. For instance, if a neuron group has 500 neurons and 128 threads are available, the number of buffer groups will be: 500/128=3.91 which is rounded up to 4. in code: (int) ceil(1.0f * groupConfigs[netId][lGrpId].numN / bufSize); For that neuron group, assuming it is the 1st neuron group in the sim, it will have buffer groups: firstNid | lastId | grpId | bufferId 0, 127, 0, 0 128, 255, 0, 1 255, 383, 0, 2 384, 499, 0, 3
In the earlier mentioned loop, nid (first neuron id in buffer group), lastId, and grpId are iterated through.
The error that is occurring is when IS_REGULAR_NEURON() attempts to identify if the nid input to the function is a correct neuron index, in the sense that it represents a real neuron index and not a false neuron index. The function IS_REGULAR_NEURON() does this by checking if the nid is less than the total number of “numNReg + numNPois” which are all regular and poisson neurons. However, what IS_REGULAR_NEURON() fails to take into account is the current buffer’s lastId.
Let’s say there are 2 neuron groups with 2 neurons in each of them. The first group would have neuron id 0 and 1 and the second group has neuron ids 2 and 3. For the first buffer position in the earlier mentioned loop, neuron ids 0, 1, 2 and 3 will be passed to the IS_REGULAR_NEURON() function. However, only neurons 0 and 1 are real neuron indices because in the first buffer the lastId is 2 meaning +2 from the firstNid of 0. IS_REGULAR_NEURON() finds that in the first buffer nids 0-3 are less than the total number of neurons in the sim and therefore reports them as real indices. What is missing here is that IS_REGULAR_NEURON() should evaluate if the nids are within firstNid to lastId of the current buffer. That would eliminate nids 2 and 3 from being reported as real neurons in the first buffer processed in the loop.
Subsequently, in the STP processing for example, nids 0 and 1 are processed twice erroneously when applying the “*= configs.dAMPA” operation. This decays the AMPA synapse current signal too fast. This same IS_REGULAR_NEURON() loop error possibly affects other operations in the GPU code as well. This error is possibly limited to GPU code because of how the threads are processed in buffers for this code. If the CPU code does similar processing it may also be an error there.