 	// int  nid  = GET_FIRING_TABLE_NID(firingId);
 //	int    post_grpId;		// STP uses pre_grpId, STDP used post_grpId...
 //	findGrpId_GPU(nid, post_grpId);
-	int post_grpId = gpuPtrs.grpIds[nid];
+	int post_grpId = runtimeDataGPU.grpIds[nid];
 
 	if(post_grpId == -1)
 		return CURRENT_UPDATE_ERROR4;
 
 	// Got one spike from dopaminergic neuron, increase dopamine concentration in the target area
-	if (gpuGrpInfo[pre_grpId].Type & TARGET_DA) {
-		atomicAdd(&(gpuPtrs.grpDA[post_grpId]), 0.04f);
+	if (groupConfigGPU[pre_grpId].Type & TARGET_DA) {
+		atomicAdd(&(runtimeDataGPU.grpDA[post_grpId]), 0.04f);
 	}
 
 	setFiringBitSynapses(nid, syn_id);
 
-	gpuPtrs.synSpikeTime[pos_ns] = simTime;		  //uncoalesced access
+	runtimeDataGPU.synSpikeTime[pos_ns] = simTime;		  //uncoalesced access
 
 	// STDP calculation: the post-synaptic neuron fires before the arrival of pre-synaptic neuron's spike
-	if (gpuGrpInfo[post_grpId].WithSTDP && !gpuNetInfo.sim_in_testing)  {
-		int stdp_tDiff = simTime-gpuPtrs.lastSpikeTime[nid];
+	if (groupConfigGPU[post_grpId].WithSTDP && !networkConfigGPU.sim_in_testing)  {
+		int stdp_tDiff = simTime-runtimeDataGPU.lastSpikeTime[nid];
 		if (stdp_tDiff >= 0) {
-			if (gpuGrpInfo[post_grpId].WithESTDP) {
+			if (groupConfigGPU[post_grpId].WithESTDP) {
 				// Handle E-STDP curves
-				switch (gpuGrpInfo[post_grpId].WithESTDPcurve) {
+				switch (groupConfigGPU[post_grpId].WithESTDPcurve) {
 				case EXP_CURVE: // exponential curve
 				case TIMING_BASED_CURVE: // sc curve
-					if (stdp_tDiff * gpuGrpInfo[post_grpId].TAU_MINUS_INV_EXC < 25.0f)
-						gpuPtrs.wtChange[pos_ns] += STDP( stdp_tDiff, gpuGrpInfo[post_grpId].ALPHA_MINUS_EXC, gpuGrpInfo[post_grpId].TAU_MINUS_INV_EXC); // uncoalesced access
+					if (stdp_tDiff * groupConfigGPU[post_grpId].TAU_MINUS_INV_EXC < 25.0f)
+						runtimeDataGPU.wtChange[pos_ns] += STDP( stdp_tDiff, groupConfigGPU[post_grpId].ALPHA_MINUS_EXC, groupConfigGPU[post_grpId].TAU_MINUS_INV_EXC); // uncoalesced access
 					break;
 				default:
 					break;
 				}
 			}
-			if (gpuGrpInfo[post_grpId].WithISTDP) {
+			if (groupConfigGPU[post_grpId].WithISTDP) {
 				// Handle I-STDP curves
-				switch (gpuGrpInfo[post_grpId].WithISTDPcurve) {
+				switch (groupConfigGPU[post_grpId].WithISTDPcurve) {
 				case EXP_CURVE: // exponential curve
-					if ((stdp_tDiff * gpuGrpInfo[post_grpId].TAU_MINUS_INV_INB) < 25.0f) { // LTD of inhibitory syanpse, which increase synapse weight
-						gpuPtrs.wtChange[pos_ns] -= STDP(stdp_tDiff, gpuGrpInfo[post_grpId].ALPHA_MINUS_INB, gpuGrpInfo[post_grpId].TAU_MINUS_INV_INB);
+					if ((stdp_tDiff * groupConfigGPU[post_grpId].TAU_MINUS_INV_INB) < 25.0f) { // LTD of inhibitory syanpse, which increase synapse weight
+						runtimeDataGPU.wtChange[pos_ns] -= STDP(stdp_tDiff, groupConfigGPU[post_grpId].ALPHA_MINUS_INB, groupConfigGPU[post_grpId].TAU_MINUS_INV_INB);
 					}
 					break;
 				case PULSE_CURVE: // pulse curve
-					if (stdp_tDiff <= gpuGrpInfo[post_grpId].LAMBDA) { // LTP of inhibitory synapse, which decreases synapse weight
-						gpuPtrs.wtChange[pos_ns] -= gpuGrpInfo[post_grpId].BETA_LTP;
-					} else if (stdp_tDiff <= gpuGrpInfo[post_grpId].DELTA) { // LTD of inhibitory syanpse, which increase synapse weight
-						gpuPtrs.wtChange[pos_ns] -= gpuGrpInfo[post_grpId].BETA_LTD;
+					if (stdp_tDiff <= groupConfigGPU[post_grpId].LAMBDA) { // LTP of inhibitory synapse, which decreases synapse weight
+						runtimeDataGPU.wtChange[pos_ns] -= groupConfigGPU[post_grpId].BETA_LTP;
+					} else if (stdp_tDiff <= groupConfigGPU[post_grpId].DELTA) { // LTD of inhibitory syanpse, which increase synapse weight
+						runtimeDataGPU.wtChange[pos_ns] -= groupConfigGPU[post_grpId].BETA_LTD;
 					}
 					break;
 				default:


	// reset various counters for the firing information
	if((blockIdx.x==0)&&(threadIdx.x==0)) {
		timingTableD1[gpuNetInfo.maxDelay]  = 0;
		spikeCountD2	+= secD2fireCnt;
		spikeCountD1	+= secD1fireCnt;
		secD2fireCnt	= timingTableD2[gpuNetInfo.maxDelay];
		secD2fireCntTest	= timingTableD2[gpuNetInfo.maxDelay];
		secD1fireCnt	= 0;
		secD1fireCntTest = 0;
	}

			// find the neuron nid and also delay information from fPos
			if((fPos>=0)&&(fPos < secD1fireCnt)) {
				atomicAdd((int*)&sh_NeuronCnt,1);
				int val  = runtimeDataGPU.firingTableD1[fPos];
				int nid  = GET_FIRING_TABLE_NID(val);
				int tPos = (networkConfigGPU.maxDelay+1)*nid;
				sh_firingId[threadIdx.x] 	 	 = val;
				sh_neuronOffsetTable[threadIdx.x]= runtimeDataGPU.cumulativePost[nid];
				sh_delayLength[threadIdx.x]      = runtimeDataGPU.postDelayInfo[tPos].delay_length;
				sh_delayIndexStart[threadIdx.x]  = runtimeDataGPU.postDelayInfo[tPos].delay_index_start;
			}
		}

