Exemple #1
0
    def compute_G(self, Dswfilename, lamb=0.0):
        """
        compute G matrix using weighting between RFs
        Dswfilename: generated by VTDM_prep
        lamb: smoothing parameter \lambda
        
        """

        Dsw = read_file(Dswfilename)
        d_Dsw = parray.to_gpu(Dsw)
        del Dsw

        #norm_func = get_put_norm_kernel(d_Dsw.dtype)
        #launch_kernel(norm_func, (256, 1, 1), (d_Dsw.shape[0],1), [d_Dsw, self.d_norm, d_Dsw.ld])

        self.d_G = parray.empty((self.size, self.size), self.dtype)

        G_func = get_G_kernel(self.dtype, d_Dsw.dtype)
        launch_kernel(G_func, (256, 1, 1), (self.d_G.shape[0], 1), [
            self.d_G, self.d_G.ld, self.d_tk1, self.d_tk2, self.Wt, self.Mt,
            d_Dsw, d_Dsw.ld, self.d_neuron_ind
        ],
                      timed="G matrix")

        if lamb != 0:
            lamb_func = get_diag_add_kernel(self.dtype)
            launch_kernel(
                lamb_func, (256, 1, 1),
                (6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, 1), [
                    self.d_G, self.d_G.ld, self.d_G.shape[0],
                    self.dtype.type(lamb)
                ])
Exemple #2
0
 def reconstruct(self, dirichfilename, time_frame, dt):
     """
     reconstruct video from c = (G)^{+}q
     
     dirichfilename: generated by either VTDM_prep or VTDM_prepb
     time_frame: a tuple or list of 2, in format [start_time, end_time]
     dt: interval between two consecutive frames in reconstruction
     
     Important Note:
     assumes the solution c is store in self.q
     
     """
 
     t = np.arange(time_frame[0], time_frame[1], dt)
     d_t = parray.to_gpu(t)
     
     dirich = read_file(dirichfilename)
     d_dirich = parray.to_gpu(dirich)
     del dirich
     
     rec_fun = get_reconstruct_kernel(d_dirich.dtype, self.d_q.dtype)
     u_rec = parray.empty((d_t.size, d_dirich.shape[1], d_dirich.shape[2]), np.float64)
     launch_kernel(rec_fun, (128,1,1), ((d_dirich.shape[1]*d_dirich.shape[2]-1) / 128+1, d_t.size), [u_rec, u_rec.ld, d_dirich, d_dirich.ld, self.d_tk1, self.d_tk2, self.d_q, d_t, self.d_neuron_ind, self.d_norm, self.Mt, self.Wt/self.Mt, self.size])
     
     return u_rec
Exemple #3
0
    def reconstruct(self, dirichfilename, time_frame, dt):
        """
        reconstruct video from c = (G)^{+}q
        
        dirichfilename: generated by either VTDM_prep or VTDM_prepb
        time_frame: a tuple or list of 2, in format [start_time, end_time]
        dt: interval between two consecutive frames in reconstruction
        
        Important Note:
        assumes the solution c is store in self.q
        
        """

        t = np.arange(time_frame[0], time_frame[1], dt)
        d_t = parray.to_gpu(t)

        dirich = read_file(dirichfilename)
        d_dirich = parray.to_gpu(dirich)
        del dirich

        rec_fun = get_reconstruct_kernel(d_dirich.dtype, self.d_q.dtype)
        u_rec = parray.empty((d_t.size, d_dirich.shape[1], d_dirich.shape[2]),
                             np.float64)
        launch_kernel(
            rec_fun, (128, 1, 1),
            ((d_dirich.shape[1] * d_dirich.shape[2] - 1) / 128 + 1, d_t.size),
            [
                u_rec, u_rec.ld, d_dirich, d_dirich.ld, self.d_tk1, self.d_tk2,
                self.d_q, d_t, self.d_neuron_ind, self.d_norm, self.Mt,
                self.Wt / self.Mt, self.size
            ])

        return u_rec
Exemple #4
0
 def compute_Dsw(self, d_Ds, Mx, My, h_norm):
     """
     Compute the weighting matrix of the "correlation" between each two RFs
     
     Parameters
     ----------
     d_Ds : PitchArray 
         containing dirichlet coefficient most possibly created by compute_Ds
     Mx : integer
         order in the x dimension
     My : integer
         order in the y dimension
     
     Returns
     -------
     PitchArray with shape (num_neurons, num_neurons)
     """
     
     if self.dtype == np.complex128:
         gemm = cublasZgemm
     else:
         gemm = cublasCgemm
     
     d_weight = parray.empty((self.num_neurons, self.num_neurons), self.dtype)
     
     handle = la.cublashandle()
     
     gemm(handle.handle, 'c', 'n', self.num_neurons, self.num_neurons, (2*Mx+1)*(2*My+1), 1.0, d_Ds.gpudata, d_Ds.ld, d_Ds.gpudata, d_Ds.ld, 0, d_weight.gpudata, d_weight.ld);
     d_Dsw = d_weight.real()
     
     norm_func = get_put_norm_kernel(d_Dsw.dtype)
     launch_kernel(norm_func, (256, 1, 1), (d_Dsw.shape[0],1), [d_Dsw, parray.to_gpu(h_norm.astype(np.float64)), d_Dsw.ld])
     
     
     return d_Dsw
Exemple #5
0
    def compute_q(self):
        """ compute q """
        self.d_q = parray.empty((self.size, 1), self.dtype)

        q_func = get_compute_q_kernel(self.dtype)
        launch_kernel(
            q_func, (256, 1, 1),
            (6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, 1), [
                self.d_q, self.d_tk1, self.d_tk2, self.d_neuron_ind,
                self.d_kappa, self.d_delta, self.d_bias, self.d_norm, self.size
            ])
Exemple #6
0
 def compute_Ds(self, Mx, My):
     """
     Parameters
     ----------
     Mx : integer
         Order in the x dimension
     My : integer
         Order in the y dimension
     
     Returns
     -------
     The dirichlet coefficients of all gabor filters with order Mx, My
     in the format of PitchArray with shape (num_neurons, 2*Mx+1, 2*My+1)
     """
     
     import scikits.cuda.cufft as cufft
     d_Ds = parray.empty((self.num_neurons, 2*My+1, 2*Mx+1), self.dtype)
     ONE_TIME_FILTER = min(1024**3 / (self.Px * self.Py * d_Ds.dtype.itemsize) / 2, self.num_neurons)
     
     n = np.asarray((self.Py, self.Px) ,np.int32)
     
     if self.dtype == np.complex128:
         plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_Z2Z, ONE_TIME_FILTER)
         fftfunc = cufft.cufftExecZ2Z
     else:
         plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_C2C, ONE_TIME_FILTER)
         fftfunc = cufft.cufftExecC2C
     
     fft2Dsfun = get_fft2Ds_kernel(dtype = self.dtype)
     
     
     for i in range(0, self.num_neurons, ONE_TIME_FILTER):
         N_filters = min(ONE_TIME_FILTER, self.num_neurons - i)
         self.generate_visual_receptive_fields(startbias = i, N_filters = N_filters)
         
         
         if N_filters < ONE_TIME_FILTER:
             cufft.cufftDestroy(plan)
             if self.dtype == np.complex128:
                 plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_Z2Z, N_filters)
                 
             else:
                 plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_C2C, N_filters)
             
         
         
         #be careful with the side-by-side constraint
         fftfunc(plan, int(self.filters.gpudata), int(self.filters.gpudata), cufft.CUFFT_FORWARD)
         
         launch_kernel(fft2Dsfun, (256, 1, 1), (Mx*2+1, My * 2+1), [[d_Ds, i * d_Ds.ld], self.filters, Mx, My, self.Px, self.Py, N_filters, d_Ds.ld, self.dx*self.dy]);
     
     cufft.cufftDestroy(plan)
     
     return d_Ds
Exemple #7
0
    def compute_dirich_space(self, d_Ds, Mx, My, Px, Py, Sx, Sy, Wx, Wy, x_start = None, y_start = None):
        """
        Compute the spatial reconstruction functions
        
        Parameters
        ----------

        d_Ds : PitchArray 
             containing dirichlet coefficient most possibly created by compute_Ds
        Mx : integer
             Order in the x dimension
        My : integer
             Order in the y dimension
        Px : integer
             number of pixels in reconstruction functions in the x dimension 
        Py : integer
             number of pixels in reconstruction functions in the y dimension
        Sx : float
             spatial domain in degree of reconstruction functions
             in x direction
        Sy : float 
             spatial domain in degree of reconstruction functions
             in y direction
        Wx : float
             spatial bandwidth in x direction 
        Wy : float
             spatial bandwidth in y direction
        x_start : float 
             indicating the starting degree in x direction
        y_start : float
             indicating the starting degree in y direction

        output: PitchArray with shape (num_neurons, Px, Py)
        """
        
        if self.dtype == np.complex128:
            typef = np.dtype(np.float64)
        else:
            typef = np.dtype(np.float32)
                
        dirich_fun = get_dirich_space_kernel(self.dtype, typef)
                
        d_dirich = parray.empty((self.num_neurons, Py, Px),typef)
            
        if x_start is None:
            x_start = - Sx/ 2

        if y_start is None:
            y_start = - Sy/2
		
        BLOCKSIZE = 16
        launch_kernel(dirich_fun,(BLOCKSIZE, BLOCKSIZE, 1), (((Px-1) / BLOCKSIZE+1) * ((Py-1) / BLOCKSIZE+1), self.num_neurons), [d_dirich, d_dirich.ld, d_Ds, d_Ds.ld, Px, Py, Mx, My, Sx, Sy, x_start, y_start, Wx / Mx, Wy / My], shared = d_Ds.dtype.itemsize * (2*Mx+1), timed = "dirich")
		
        return d_dirich
Exemple #8
0
    def generate_visual_receptive_fields(self, startbias = 0, N_filters = None, x_start = None, y_start = None):
        """
        Generate a batch of centre surround filters from parameters set in self
        
        Parameters
        ----------
        start_bias : integer, optional
            start the the (start_bias)th filter
        N_filters : integer, optional 
            generate N_filters filters
        x_start : float
            indicating the starting degree in x direction
        y_start : float 
            indicating the starting degree in y direction

        
        """
        if N_filters is None:
            N_filters = self.num_neurons - startbias
        
        
        try:
            if N_filters > self.filters.shape[0]:
                del self.filters
                self.filters = parray.empty((N_filters, self.Pyall, self.Pxall), self.dtype)
        except:
            self.filters = parray.empty((N_filters, self.Pyall, self.Pxall), self.dtype)
        

        if x_start is None:
            x_start = - self.Sxall/ 2

        if y_start is None:
            y_start = -self.Syall/2
        
        BLOCK_SIZE = 16
        
        launch_kernel(self.func, (BLOCK_SIZE, BLOCK_SIZE, 1), (((self.Pxall-1)/BLOCK_SIZE+1) * ((self.Pyall-1)/BLOCK_SIZE+1), int(N_filters)), [self.filters, self.filters.ld, [self.d_alpha, startbias], [self.d_x0, startbias], [self.d_y0, startbias], self.Pxall, self.Pyall, self.Sxall, self.Syall, x_start, y_start, self.sigma_center**2, self.sigma_surround**2])
Exemple #9
0
    def compute_G(self, Dswfilename, lamb=0.0):
        """
        compute G matrix using weighting between RFs
        Dswfilename: generated by VTDM_prep
        lamb: smoothing parameter \lambda
        
        """
        
        
        Dsw = read_file(Dswfilename)
        d_Dsw = parray.to_gpu(Dsw)
        del Dsw

        #norm_func = get_put_norm_kernel(d_Dsw.dtype)
        #launch_kernel(norm_func, (256, 1, 1), (d_Dsw.shape[0],1), [d_Dsw, self.d_norm, d_Dsw.ld])

        self.d_G = parray.empty((self.size, self.size), self.dtype)

        G_func = get_G_kernel(self.dtype, d_Dsw.dtype)
        launch_kernel(G_func, (256, 1, 1), (self.d_G.shape[0], 1), [self.d_G, self.d_G.ld, self.d_tk1, self.d_tk2, self.Wt, self.Mt, d_Dsw, d_Dsw.ld, self.d_neuron_ind], timed = "G matrix")

        if lamb != 0:
            lamb_func = get_diag_add_kernel(self.dtype)
            launch_kernel(lamb_func, (256,1,1), (6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, 1), [self.d_G, self.d_G.ld, self.d_G.shape[0], self.dtype.type(lamb)])
Exemple #10
0
    def compute_Gb(self, Dsfilename, lamb=0.0):
        """
        compute G matrix using dirichlet coefficients
        Dsfilename: generated by VTDM_prepb
        lamb: smoothing parameter \lambda
        
        """
        handle = la.cublashandle()
        import tables
        h5file = tables.openFile(Dsfilename)
        Ds = h5file.root.real.read()

        d_Ds = parray.to_gpu(Ds.reshape((Ds.shape[0], -1)))
        del Ds

        d_Dsw = parray.empty((d_Ds.shape[0], d_Ds.shape[0]), d_Ds.dtype)
        if d_Ds.dtype == np.float64:
            from scikits.cuda.cublas import cublasDgemm
            gemm = cublasDgemm
        else:
            from scikits.cuda.cublas import cublasSgemm
            gemm = cublasSgemm
        gemm(handle.handle, 't', 'n', d_Dsw.shape[0], d_Dsw.shape[0],
             d_Ds.shape[1], 1.0, d_Ds.gpudata, d_Ds.ld, d_Ds.gpudata, d_Ds.ld,
             0.0, d_Dsw.gpudata, d_Dsw.ld)
        Ds = h5file.root.imag.read()
        d_Ds.set(Ds)
        gemm(handle.handle, 't', 'n', d_Dsw.shape[0], d_Dsw.shape[0],
             d_Ds.shape[1], 1.0, d_Ds.gpudata, d_Ds.ld, d_Ds.gpudata, d_Ds.ld,
             1.0, d_Dsw.gpudata, d_Dsw.ld)
        del Ds
        h5file.close()

        norm_func = get_put_norm_kernel(d_Dsw.dtype)
        launch_kernel(norm_func, (256, 1, 1), (d_Dsw.shape[0], 1),
                      [d_Dsw, self.d_norm, d_Dsw.ld])

        self.d_G = parray.empty((self.size, self.size), self.dtype)

        G_func = get_G_kernel(self.dtype, d_Dsw.dtype)
        launch_kernel(G_func, (256, 1, 1), (self.d_G.shape[0], 1), [
            self.d_G, self.d_G.ld, self.d_tk1, self.d_tk2, self.Wt, self.Mt,
            d_Dsw, d_Dsw.ld, self.d_neuron_ind
        ],
                      timed="G matrix")

        if lamb != 0:
            lamb_func = get_diag_add_kernel(self.dtype)
            launch_kernel(
                lamb_func, (256, 1, 1),
                (6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, 1), [
                    self.d_G, self.d_G.ld, self.d_G.shape[0],
                    self.dtype.type(lamb)
                ])
Exemple #11
0
    def compute_Gb(self, Dsfilename, lamb=0.0):
        """
        compute G matrix using dirichlet coefficients
        Dsfilename: generated by VTDM_prepb
        lamb: smoothing parameter \lambda
        
        """
        handle = la.cublashandle()
        import tables
        h5file = tables.openFile(Dsfilename)
        Ds = h5file.root.real.read()
        
        d_Ds = parray.to_gpu(Ds.reshape((Ds.shape[0],-1)))
        del Ds
 
        d_Dsw = parray.empty((d_Ds.shape[0], d_Ds.shape[0]), d_Ds.dtype)
        if d_Ds.dtype == np.float64:
            from scikits.cuda.cublas import cublasDgemm
            gemm = cublasDgemm
        else:
            from scikits.cuda.cublas import cublasSgemm
            gemm = cublasSgemm
        gemm(handle.handle, 't', 'n', d_Dsw.shape[0], d_Dsw.shape[0], d_Ds.shape[1], 1.0, d_Ds.gpudata, d_Ds.ld, d_Ds.gpudata, d_Ds.ld, 0.0, d_Dsw.gpudata, d_Dsw.ld)
        Ds = h5file.root.imag.read()
        d_Ds.set(Ds)
        gemm(handle.handle, 't', 'n', d_Dsw.shape[0], d_Dsw.shape[0], d_Ds.shape[1], 1.0, d_Ds.gpudata, d_Ds.ld, d_Ds.gpudata, d_Ds.ld, 1.0, d_Dsw.gpudata, d_Dsw.ld)        
        del Ds
        h5file.close()
        
        norm_func = get_put_norm_kernel(d_Dsw.dtype)
        launch_kernel(norm_func, (256, 1, 1), (d_Dsw.shape[0],1), [d_Dsw, self.d_norm, d_Dsw.ld])

        self.d_G = parray.empty((self.size, self.size), self.dtype)

        G_func = get_G_kernel(self.dtype, d_Dsw.dtype)
        launch_kernel(G_func, (256, 1, 1), (self.d_G.shape[0], 1), [self.d_G, self.d_G.ld, self.d_tk1, self.d_tk2, self.Wt, self.Mt, d_Dsw, d_Dsw.ld, self.d_neuron_ind], timed = "G matrix")

        if lamb != 0:
            lamb_func = get_diag_add_kernel(self.dtype)
            launch_kernel(lamb_func, (256,1,1), (6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, 1), [self.d_G, self.d_G.ld, self.d_G.shape[0], self.dtype.type(lamb)])
Exemple #12
0
 def compute_q(self ):
     """ compute q """
     self.d_q = parray.empty((self.size,1), self.dtype)
     
     q_func = get_compute_q_kernel(self.dtype)
     launch_kernel(q_func, (256, 1, 1), (6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, 1), [self.d_q, self.d_tk1, self.d_tk2, self.d_neuron_ind, self.d_kappa, self.d_delta, self.d_bias, self.d_norm, self.size])
Exemple #13
0
 def encode(self, neural_inputs, startbias = 0, avg_rate = 0.1):
     """
         Encode with IAFs with random thresholds
         
         Parameters
         ----------
         neural_inputs : PitchArray 
              PitchArray of shape (num_samples, num_neurons) containing inputs to all neurons
         startbias : integer 
              the neuron index corresponding to first column of neural_inputs
         avg_rate : float 
              average spiking rate assumed for neurons, will allocate memory num_samples/avg_rate for each neuron for storing spikes
         
         Returns
         -------
         spikes : ndarray of self.dtype 
              stores the spikes for one neuron after another
         spike_count : ndarray of int32 of size num_neurons 
              indicates the number of spikes generated by each neuron
         
         Notes
         -----
         spikes for neuron j can be accessed by
         ::
         
              cum_count = np.concatenate((np.zeros(1,np.int32),np.cumsum(spike_count)))
              tk = spikes[cum_count[j]:cum_count[j+1]]
         
         """
     
     neuron_per_block=64
     
     if self.num_neurons != neural_inputs.shape[1]:
         raise ValueError("input size should match number of neurons")
     
     Ntimesteps = neural_inputs.shape[0]
     
     d_spikecount = parray.empty((1, self.num_neurons), np.int32)
     
     randnum = np.random.normal(size = ( int(np.ceil(Ntimesteps / avg_rate)), self.num_neurons)).astype(self.dtype)
     #d_spike = parray.empty( ( int(np.ceil(Ntimesteps / avg_rate)), self.num_neurons), self.dtype) 
     d_spike = parray.to_gpu(randnum)
     
     
     if neural_inputs.__class__ is np.ndarray:
         d_neural_inputs = parray.to_gpu(neural_inputs)
     else:
         d_neural_inputs = neural_inputs
     
     
     launch_kernel(self.func, (neuron_per_block, 1, 1), (int(np.ceil(np.float64(self.num_neurons) / neuron_per_block)), 1), \
                   [neural_inputs, neural_inputs.ld, self.num_neurons, Ntimesteps, d_spike, d_spike.ld, [self.d_v0, startbias], \
                    [self.d_kappa, startbias], [self.d_bias, startbias], [self.d_delta, startbias], [self.d_time_count, startbias],\
                    d_spikecount, int(np.ceil(Ntimesteps / avg_rate)), self.dt, [self.d_delta_value,startbias], [self.d_sigma,startbias]], shared = self.dtype.itemsize * neuron_per_block)
     
     
     spike_count = d_spikecount.get()
     spike_count.resize((self.num_neurons,))
     
     if spike_count.max() >= np.ceil(Ntimesteps / avg_rate):
         raise ValueError("number of spikes exceeded the limit of buffer")
     
     spike = rearrange_spikes(d_spike, spike_count, self.num_neurons)
     
     return spike, spike_count
Exemple #14
0
 def compute_dirich_space_fft(self, d_Ds, Mx, My, Px, Py, Sx, Sy, Wx, Wy):
     import scikits.cuda.cufft as cufft
     
     dx = Sx / Px
     dy = Sy / Py
     
     Px1 = int(np.round(self.Sx / dx))
     Py1 = int(np.round(self.Sy / dy))
     
     
     if self.dtype == np.complex128:
         typef = np.dtype(np.float64)
     else:
         typef = np.dtype(np.float32)
     
     d_dirich = parray.empty((self.num_neurons, Py, Px),typef)
     
     freemem,totalmem = cuda.mem_get_info()
     
     ONE_TIME_FILTER = int(min(freemem / (Px1 * Py1 * d_Ds.dtype.itemsize) / 4, self.num_neurons))
     
     
     n = np.asarray((Py1, Px1) ,np.int32)
     
     
     if self.dtype == np.complex128:
         plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_Z2Z, ONE_TIME_FILTER)
         fftfunc = cufft.cufftExecZ2Z
     else:
         plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_C2C, ONE_TIME_FILTER)
         fftfunc = cufft.cufftExecC2C
     
     
     Ds2fftfun = get_Ds2fft_kernel(self.dtype)
     d_filter_complex = parray.empty((ONE_TIME_FILTER, Px1*Py1), self.dtype)
     filter2recfun = get_filter2rec_kernel(self.dtype)
     
     for i in range(0, self.num_neurons, ONE_TIME_FILTER):
         N_filters = min(ONE_TIME_FILTER, self.num_neurons - i)
         d_filter_complex.fill(0)
         
         launch_kernel(Ds2fftfun, (256,1,1), (Mx*2+1, My*2+1), [[d_Ds,i * d_Ds.ld], d_Ds.ld, d_filter_complex, d_filter_complex.ld, Mx, My, Px1, Py1, N_filters])
         
         if N_filters < ONE_TIME_FILTER:
             cufft.cufftDestroy(plan)
             if self.dtype == np.complex128:
                 plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_Z2Z, N_filters)
             
             else:
                 plan = cufft.cufftPlanMany(2, n.ctypes.data, None, 1, 0, None, 1, 0, cufft.CUFFT_C2C, N_filters)
         
         
         
         #be careful with the side-by-side constraint
         fftfunc(plan, int(d_filter_complex.gpudata), int(d_filter_complex.gpudata), cufft.CUFFT_INVERSE)
         
         
         BLOCK_SIZE = 16
         launch_kernel(filter2recfun, (BLOCK_SIZE,BLOCK_SIZE,1), (((Px-1) / BLOCK_SIZE + 1)* ((Py-1) / BLOCK_SIZE+1), N_filters), [[d_dirich,i * d_dirich.ld],d_dirich.ld, d_filter_complex, d_filter_complex.ld, N_filters, Px, Py, Px1, Py1])
     
     cufft.cufftDestroy(plan)
     
     return d_dirich