 	}
 }
 
-////////////////////////////////////////////////////////////////////////////////////////////////
-/// Device local function:      gpu_newFireUpdate                                           ///
-/// Description: 1. Copy neuron id from local table to global firing table.                 ///
-///		 2. Reset all neuron properties of neuron id in local table 		    			///
-//                                                                                          ///
-/// fireTablePtr:  local shared memory firing table with neuron ids of fired neuron         ///
-/// fireCntD2:      number of excitatory neurons in local table that has fired               ///
-/// fireCntD1:      number of inhibitory neurons in local table that has fired               ///
-/// simTime:      current global time step..stored as neuron firing time  entry             ///
-////////////////////////////////////////////////////////////////////////////////////////////////
+/*
+ * \brief 1. Copy neuron id from local table to global firing table. 2. Reset all neuron properties of neuron id in local table
+ *
+ *
+ * \param[in] fireTablePtr the local shared memory firing table with neuron ids of fired neuron
+ * \param[in] fireCntD2 the number of neurons in local table that has fired with group's max delay == 1
+ * \param[in] fireCntD1 the number of neurons in local table that has fired with group's max delay > 1
+ * \param[in] simTime the current time step, stored as neuron firing time  entry
+ */
 __device__ void updateFiringCounter(volatile unsigned int& fireCnt, volatile unsigned int& fireCntD1, volatile unsigned int& cntD2, volatile unsigned int& cntD1, volatile int&  blkErrCode)
 {
-	int fireCntD2 = fireCnt-fireCntD1;
+	int fireCntD2 = fireCnt - fireCntD1;
 
 	cntD2 = atomicAdd(&secD2fireCntTest, fireCntD2);
 	cntD1 = atomicAdd(&secD1fireCntTest, fireCntD1);

	return;
}

///////////////////////////////////////////////////////////////////////////
/// Device local function:      gpu_updateLTP
/// Description:                Computes the STDP update values for each of fired
///                             neurons stored in the local firing table.
///
/// fireTablePtr:  local shared memory firing table with neuron ids of fired neuron
/// fireCnt:      number of fired neurons in local firing table
/// simTime:     current global time step..stored as neuron firing time  entry
///////////////////////////////////////////////////////////////////////////
// synaptic grouping for LTP Calculation
#define		LTP_GROUPING_SZ     16
__device__ void gpu_updateLTP(	int*     		fireTablePtr,
				short int*  		fireGrpId,
				volatile unsigned int&   fireCnt,
				int&      		simTime)
{
	for(int pos=threadIdx.x/LTP_GROUPING_SZ; pos < fireCnt; pos += (blockDim.x/LTP_GROUPING_SZ))  {
		// each neuron has two variable pre and pre_exc
		// pre: number of pre-neuron
		// pre_exc: number of neuron had has plastic connections
		short int grpId = fireGrpId[pos];
		// STDP calculation: the post-synaptic neron fires after the arrival of pre-synaptic neuron's spike
		if (groupConfigGPU[grpId].WithSTDP) { // MDR, FIXME this probably will cause more thread divergence than need be...
			int  nid   = fireTablePtr[pos];


			__syncthreads();

			// the local firing table is full. dump the local firing table to the global firing table before proceeding
			if (fireCntTest >= (FIRE_CHUNK_CNT)) {

				// clear the table and update...
				int retCode = updateNewFirings(fireTable,  fireGrpId, fireCnt, fireCntD1, simTime);
				if (retCode != 0) return;
				// update based on stdp rule
				// KILLME !!! if (simTime > 0))
