#include "definitionsInternal.h"
#include "supportCode.h"

extern "C" __global__ void preNeuronResetKernel() {
    unsigned int id = 32 * blockIdx.x + threadIdx.x;
    if(id == 0) {
        dd_glbSpkCntE[0] = 0;
    }
    else if(id == 1) {
        dd_glbSpkCntI[0] = 0;
    }
}
extern "C" __global__ void updateNeuronsKernel(float t)
 {
    const unsigned int id = 128 * blockIdx.x + threadIdx.x;
    __shared__ volatile unsigned int shSpk[128];
    __shared__ volatile unsigned int shPosSpk;
    __shared__ volatile unsigned int shSpkCount;
    if (threadIdx.x == 0); {
        shSpkCount = 0;
    }

    __syncthreads();
    // E
    if(id < 3200) {

        if(id < 3200) {
            scalar lV = dd_VE[id];
            scalar lRefracTime = dd_RefracTimeE[id];

            float Isyn = 0;
            // pull inSyn values in a coalesced access
            float linSynIE = dd_inSynIE[id];
            Isyn += (9.51625819640404824e-01f) * linSynIE;
            // pull inSyn values in a coalesced access
            float linSynEE = dd_inSynEE[id];
            Isyn += (9.06346234610090895e-01f) * linSynEE;
            // test whether spike condition was fulfilled previously
            const bool oldSpike= (lRefracTime <= 0.0f && lV >= (-5.00000000000000000e+01f));
            // calculate membrane potential
            if (lRefracTime <= 0.0f)
            {
              scalar alpha = ((Isyn + (0.00000000000000000e+00f)) * (2.00000000000000000e+01f)) + (-4.90000000000000000e+01f);
              lV = alpha - ((9.51229424500714016e-01f) * (alpha - lV));
            }
            else
            {
              lRefracTime -= DT;
            }

            // test for and register a true spike
            if ((lRefracTime <= 0.0f && lV >= (-5.00000000000000000e+01f)) && !(oldSpike)) {
                const unsigned int spkIdx = atomicAdd((unsigned int *) &shSpkCount, 1);
                shSpk[spkIdx] = id;
                // spike reset code
                lV = (-6.00000000000000000e+01f);
                lRefracTime = (5.00000000000000000e+00f);

            }
            dd_VE[id] = lV;
            dd_RefracTimeE[id] = lRefracTime;
            // the post-synaptic dynamics
            linSynIE*=(9.04837418035959518e-01f);
            dd_inSynIE[id] = linSynIE;
            // the post-synaptic dynamics
            linSynEE*=(8.18730753077981821e-01f);
            dd_inSynEE[id] = linSynEE;
        }
        __syncthreads();
        if (threadIdx.x == 0) {
            if (shSpkCount > 0) {
                shPosSpk = atomicAdd((unsigned int *) &dd_glbSpkCntE[0], shSpkCount);
            }
        }
        __syncthreads();
        if (threadIdx.x < shSpkCount) {
            const unsigned int n = shSpk[threadIdx.x];
            dd_glbSpkE[shPosSpk + threadIdx.x] = n;
        }
    }

    // I
    if(id >= 3200 && id < 4096) {
        const unsigned int lid = id - 3200;

        if(lid < 800) {
            scalar lV = dd_VI[lid];
            scalar lRefracTime = dd_RefracTimeI[lid];

            float Isyn = 0;
            // pull inSyn values in a coalesced access
            float linSynII = dd_inSynII[lid];
            Isyn += (9.51625819640404824e-01f) * linSynII;
            // pull inSyn values in a coalesced access
            float linSynEI = dd_inSynEI[lid];
            Isyn += (9.06346234610090895e-01f) * linSynEI;
            // test whether spike condition was fulfilled previously
            const bool oldSpike= (lRefracTime <= 0.0f && lV >= (-5.00000000000000000e+01f));
            // calculate membrane potential
            if (lRefracTime <= 0.0f)
            {
              scalar alpha = ((Isyn + (0.00000000000000000e+00f)) * (2.00000000000000000e+01f)) + (-4.90000000000000000e+01f);
              lV = alpha - ((9.51229424500714016e-01f) * (alpha - lV));
            }
            else
            {
              lRefracTime -= DT;
            }

            // test for and register a true spike
            if ((lRefracTime <= 0.0f && lV >= (-5.00000000000000000e+01f)) && !(oldSpike)) {
                const unsigned int spkIdx = atomicAdd((unsigned int *) &shSpkCount, 1);
                shSpk[spkIdx] = lid;
                // spike reset code
                lV = (-6.00000000000000000e+01f);
                lRefracTime = (5.00000000000000000e+00f);

            }
            dd_VI[lid] = lV;
            dd_RefracTimeI[lid] = lRefracTime;
            // the post-synaptic dynamics
            linSynII*=(9.04837418035959518e-01f);
            dd_inSynII[lid] = linSynII;
            // the post-synaptic dynamics
            linSynEI*=(8.18730753077981821e-01f);
            dd_inSynEI[lid] = linSynEI;
        }
        __syncthreads();
        if (threadIdx.x == 0) {
            if (shSpkCount > 0) {
                shPosSpk = atomicAdd((unsigned int *) &dd_glbSpkCntI[0], shSpkCount);
            }
        }
        __syncthreads();
        if (threadIdx.x < shSpkCount) {
            const unsigned int n = shSpk[threadIdx.x];
            dd_glbSpkI[shPosSpk + threadIdx.x] = n;
        }
    }

}
void updateNeurons(float) {
     {
        const dim3 threads(32, 1);
        const dim3 grid(1, 1);
        preNeuronResetKernel<<<grid, threads>>>();
    }
     {
        cudaEventRecord(neuronUpdateStart);
        const dim3 threads(128, 1);
        const dim3 grid(32, 1);
        updateNeuronsKernel<<<grid, threads>>>(t);
        cudaEventRecord(neuronUpdateStop);
    }
}