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
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
def load_parameters(self, num_neurons = None, h_alpha = None, h_l = None, h_x0 = None, h_y0 = None, h_ab = None, KAPPA = 2.5, set = 0): """ Load gabor parameters to GPU num_neurons, h_alpha, h_l, h_x0, h_y0, h_ab must be specified together or not specified at all Parameters ---------- num_neurons : integer total number of neurons h_alpha : ndarray of float64 containing dilation parameters alpha = alpha0**(m) h_l : ndarray of float64 containing rotation parameters l (angles) h_x0 : ndarray of float64 containing translation parameters x0 h_y0 : ndarray of float64 containing translation parameters y0 h_ab : ndarray of int32 containing 0 or 1, 0 for real part, 1 for imaginary part jth gabor filter will be generated according to h_alpha[j], h_l[j], h_n[j], h_k[j] and h_ab[j] KAPPA : float spatial bandwidth set : integer 0 if self parameters has not been set, 1 if they have been set by get_gabor_parameters or other manner) """ if num_neurons is not None: self.num_neurons = num_neurons if h_alpha is None or h_l is None or h_x0 is None or h_y0 is None or h_ab is None: raise ValueError("must specify the gabor parameters") else: self.h_alpha = h_alpha self.h_l = h_l self.h_x0 = h_x0 self.h_y0 = h_y0 self.h_ab = h_ab self.KAPPA = KAPPA else: if not set: self.get_gabor_parameters(KAPPA=KAPPA) self.d_alpha = parray.to_gpu(1.0 / self.h_alpha) self.d_l = parray.to_gpu(self.h_l) self.d_x0 = parray.to_gpu(self.h_x0) self.d_y0 = parray.to_gpu(self.h_y0) self.d_ab = parray.to_gpu(self.h_ab)
def load_parameters(self, num_neurons = None, h_alpha = None, h_x0 = None, h_y0 = None, set = 0, sigma_center = 0.5, sigma_surround = 0.8): """ Load Centre Surround parameters to GPU num_neurons, h_alpha, h_l, h_x0, h_y0, h_ab must be specified together or unspecified together Parameters ---------- num_neurons : integer total number of neurons h_alpha : ndarray of float64 containing dilation parameters alpha = alpha0**(m) h_l : ndarray of float64 containing rotation parameters l (angles) h_x0 : ndarray of float64 containing translation parameters x0 h_y0 : ndarray of float64 containing translation parameters y0 h_ab : ndarray of int32 containing 0 or 1, 0 for real part, 1 for imaginary part jth gabor filter will be generated according to h_alpha[j], h_l[j], h_n[j], h_k[j] and h_ab[j] set : integer 0 if self parameters has not been set, 1 if they have been set by get_gabor_parameters or other manner) sigma_center : float, optional standard deviation of the center sigma_surround : float, optional standard deviation of the surround """ if num_neurons is not None: self.num_neurons = num_neurons if h_alpha is None or h_x0 is None or h_y0 is None: raise ValueError("must specify the gabor parameters") else: self.h_alpha = h_alpha self.h_x0 = h_x0 self.h_y0 = h_y0 self.sigma_center = float(sigma_center) self.sigma_surround = float(sigma_surround) else: if not set: self.get_cs_parameters() self.d_alpha = parray.to_gpu(1.0 / self.h_alpha) self.d_x0 = parray.to_gpu(self.h_x0) self.d_y0 = parray.to_gpu(self.h_y0)
def filter(self, V): """ Filter a video V Must set up parameters of CS RF first Parameters ---------- V : 3D ndarray, with shape (num_frames, Px, Py) Returns ------- the filtered output by the gabor filters specified in self output is a PitchArray with shape (num_neurons, num_frames), jth row of which is the output of jth gabor filter """ d_output = parray.empty((self.num_neurons, V.shape[0]), self.dtype) d_video = parray.to_gpu(V.reshape(V.shape[0], V.shape[1]*V.shape[2])) free,total = cuda.mem_get_info() self.ONE_TIME_FILTERS = (free / self.dtype.itemsize) * 3/4 / self.Pxall / self.Pyall handle = la.cublashandle() for i in np.arange(0,self.num_neurons,self.ONE_TIME_FILTERS): Nfilters = min(self.ONE_TIME_FILTERS, self.num_neurons - i) self.generate_visual_receptive_fields(startbias = i, N_filters = Nfilters) cublasDgemm(handle.handle, 't','n', V.shape[0], int(Nfilters), self.Pxall*self.Pyall, self.dx*self.dy, d_video.gpudata, d_video.ld, self.filters.gpudata, self.filters.ld, 0, int(int(d_output.gpudata)+int(d_output.ld*i*d_output.dtype.itemsize)) , d_output.ld) return d_output.T()
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) ])
def filter(self, video_input): """ Performs RF filtering on input video for all the rfs """ if len(video_input.shape) == 2: # if input has 2 dimensions assert video_input.shape[1] == self.size else: # if input has 3 dimensions assert (video_input.shape[1]*video_input.shape[2] == self.size) # rasterizing inputs video_input.resize((video_input.shape[0], self.size)) d_video = parray.to_gpu(video_input) d_output = parray.empty((self.num_neurons, video_input.shape[0]), self.dtype) free, total = cuda.mem_get_info() self.ONE_TIME_FILTERS = ((free // self.dtype.itemsize) * 3 // 4 // self.size) self.ONE_TIME_FILTERS -= self.ONE_TIME_FILTERS % 2 self.ONE_TIME_FILTERS = min(self.ONE_TIME_FILTERS, self.num_neurons) handle = la.cublashandle() for i in np.arange(0, self.num_neurons, self.ONE_TIME_FILTERS): Nfilters = min(self.ONE_TIME_FILTERS, self.num_neurons - i) self.generate_filters(startbias=i, N_filters=Nfilters) la.dot(self.filters, d_video, opb='t', C=d_output[i: i+Nfilters], handle=handle) del self.filters return d_output.T()
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
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) ])
def pre_run(self): super(LPU, self).pre_run() self._setup_connectivity() self._initialize_gpu_ds() self._init_objects() self.buffer = circular_array(self.total_gpot_neurons, self.my_num_gpot_neurons,\ self.gpot_delay_steps, self.V, \ self.total_spike_neurons, self.spike_delay_steps) if self.input_file: self.input_h5file = tables.openFile(self.input_file) self.file_pointer = 0 self.I_ext = parray.to_gpu(self.input_h5file.root.array.read(\ self.file_pointer, self.file_pointer + \ self._one_time_import)) self.file_pointer += self._one_time_import self.frame_count = 0 self.frames_in_buffer = self._one_time_import if self.output: output_file = self.output_file.rsplit('.', 1) filename = output_file[0] if (len(output_file) > 1): ext = output_file[1] else: ext = 'h5' if self.my_num_gpot_neurons > 0: self.output_gpot_file = tables.openFile(filename + \ '_gpot.' + ext , 'w') self.output_gpot_file.createEArray("/","array", \ tables.Float64Atom(), (0,self.my_num_gpot_neurons)) if self.my_num_spike_neurons > 0: self.output_spike_file = tables.openFile(filename + \ '_spike.' + ext , 'w') self.output_spike_file.createEArray("/","array", \ tables.Float64Atom(),(0,self.my_num_spike_neurons)) if self.debug: self.in_gpot_files = {} for (key, i) in self.other_lpu_map.iteritems(): num = self.num_input_gpot_neurons[i] if num > 0: self.in_gpot_files[key] = tables.openFile(filename + \ key + '_in_gpot.' + ext , 'w') self.in_gpot_files[key].createEArray("/","array", \ tables.Float64Atom(), (0,num)) self.gpot_buffer_file = tables.openFile(self.id + '_buffer.h5', 'w') self.gpot_buffer_file.createEArray("/","array", \ tables.Float64Atom(), (0,self.gpot_delay_steps, self.total_gpot_neurons))
def pre_run(self): super(LPU, self).pre_run() self._setup_connectivity() self._initialize_gpu_ds() self._init_objects() self.buffer = circular_array(self.total_gpot_neurons, self.my_num_gpot_neurons,\ self.gpot_delay_steps, self.V, \ self.total_spike_neurons, self.spike_delay_steps) if self.input_file: self.input_h5file = tables.openFile(self.input_file) self.file_pointer = 0 self.I_ext = parray.to_gpu(self.input_h5file.root.array.read(\ self.file_pointer, self.file_pointer + \ self._one_time_import)) self.file_pointer += self._one_time_import self.frame_count = 0 self.frames_in_buffer = self._one_time_import if self.output: output_file = self.output_file.rsplit('.',1) filename = output_file[0] if(len(output_file)>1): ext = output_file[1] else: ext = 'h5' if self.my_num_gpot_neurons>0: self.output_gpot_file = tables.openFile(filename + \ '_gpot.' + ext , 'w') self.output_gpot_file.createEArray("/","array", \ tables.Float64Atom(), (0,self.my_num_gpot_neurons)) if self.my_num_spike_neurons>0: self.output_spike_file = tables.openFile(filename + \ '_spike.' + ext , 'w') self.output_spike_file.createEArray("/","array", \ tables.Float64Atom(),(0,self.my_num_spike_neurons)) if self.debug: self.in_gpot_files = {} for (key,i) in self.other_lpu_map.iteritems(): num = self.num_input_gpot_neurons[i] if num>0: self.in_gpot_files[key] = tables.openFile(filename + \ key + '_in_gpot.' + ext , 'w') self.in_gpot_files[key].createEArray("/","array", \ tables.Float64Atom(), (0,num)) self.gpot_buffer_file = tables.openFile(self.id + '_buffer.h5','w') self.gpot_buffer_file.createEArray("/","array", \ tables.Float64Atom(), (0,self.gpot_delay_steps, self.total_gpot_neurons))
def load_parameters(self, h_kappa = None, h_bias = None, h_delta = None, h_sigma = None, h_time_count = None, h_v0 = None): """ Load encoding parameters to GPU h_kappa, h_bias, h_delta can be set up using default values (using None), or specified together (using ndarrays of dtype) h_time_count can be set up using default values (using None), or specified values(using ndarrays of dtype) h_v0 can be set up using default values (using None), or specified values(using ndarrays of dtype) """ if h_kappa is None: self.set_parameters() else: self.h_kappa = h_kappa self.h_delta = h_delta self.h_bias = h_bias self.h_sigma = np.zeros(h_kappa.shape) if h_time_count is None: self.set_time_count() else: self.h_time_count = h_time_count if h_v0 is None: self.set_initial_value() else: self.h_v0 = h_v0 self.d_kappa = parray.to_gpu(self.h_kappa) self.d_delta = parray.to_gpu(self.h_delta) self.d_bias = parray.to_gpu(self.h_bias) self.d_v0 = parray.to_gpu(self.h_v0) self.d_time_count = parray.to_gpu(self.h_time_count)
def __init__(self, decode_neuron, tk1, tk2, neuron_ind, kappa, delta, bias, norm, Wt, Mt, dtype=np.float64): """ class dirichlet For use to reconstruct videos encoded by VTEMs consist of spatial only receptive fields and IAF neurons parameters: decode_neuron: number of neuron used in the decoding tk1: a ndarray vector of float64 containing all the t_k's in \int_{t_k}^{t_{k+1}} tk2: a ndarray vector of float64 containing all the t_{k+1}'s in \int_{t_k}^{k+1}} neuron_ind: a ndarray vector of int32, where neuron_ind[j] is the neuron index correspond to (tk1[j],tk2[j]) kappa: a ndarray vector of float64 containing the integration constants for each neuron delta: a ndarray vector of float64 containing the threshold for each neuron bias: a ndarray vector of float64 containing the bias for each neuron norm: a ndarray vector of float64 containing the normalization factor for each neuron (in case of deterministic threshold, use a vector of ones) Wt: temporal bandwidth Mt: dirichlet order in variable t dtype: np.float64 or np.float32, determine dtype of G matrix and q vector """ self.size = tk1.size if self.size == 0: raise ValueError("no spikes to decode") self.Wt = Wt self.Mt = Mt self.dtype = np.dtype(dtype) print "number of spikes: %d" % (self.size) self.d_tk1 = parray.to_gpu(tk1) self.d_tk2 = parray.to_gpu(tk2) self.d_neuron_ind = parray.to_gpu(neuron_ind) self.d_kappa = parray.to_gpu(kappa) self.d_delta = parray.to_gpu(delta) self.d_bias = parray.to_gpu(bias) self.d_norm = parray.to_gpu(norm)
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)])
def pre_run(self): super(LPU,self).pre_run() self._setup_connectivity() self._initialize_gpu_ds() self._init_objects() self.buffer = circular_array(self.total_gpot_neurons, self.my_num_gpot_neurons,\ self.gpot_delay_steps, self.V, \ self.total_spike_neurons, self.spike_delay_steps) if self.input_file is not None: self.input_h5file = tables.openFile(self.input_file) self.one_time_import = 10 self.file_pointer = 0 self.I_ext = parray.to_gpu(self.input_h5file.root.array.read(\ self.file_pointer, self.file_pointer + \ self.one_time_import)) self.file_pointer += self.one_time_import self.frame_count = 0 if self.output: output_file = self.output_file.rsplit('.',1) filename = output_file[0] if(len(output_file)>1): ext = output_file[1] else: ext = 'h5' if self.my_num_gpot_neurons>0: self.output_gpot_file = tables.openFile(filename + \ '_gpot.' + ext , 'w') self.output_gpot_file.createEArray("/","array", \ tables.Float64Atom(), (0,self.my_num_gpot_neurons)) if self.my_num_spike_neurons>0: self.output_spike_file = tables.openFile(filename + \ '_spike.' + ext , 'w') self.output_spike_file.createEArray("/","array", \ tables.Float64Atom(),(0,self.my_num_spike_neurons))
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)])
def filter_image(self, image_input): """ Performs RF filtering on input video for all the rfs """ # video dimensions should match screen dimensions # numpy resize operation doesn,t make any checks if len(image_input.shape) == 2: # if input has 2 dimensions assert image_input.shape[1] == self.size else: # if input has 3 dimensions assert (image_input.shape[1]*image_input.shape[2] == self.size) # rasterizing inputs image_input.resize((1, self.size)) d_image = parray.to_gpu(image_input) d_output = parray.empty((self.num_neurons, image_input.shape[0]), self.dtype) free, total = cuda.mem_get_info() self.ONE_TIME_FILTERS = ((free // self.dtype.itemsize) * 3 // 4 // self.size) self.ONE_TIME_FILTERS -= self.ONE_TIME_FILTERS % 2 self.ONE_TIME_FILTERS = min(self.ONE_TIME_FILTERS, self.num_neurons) handle = la.cublashandle() for i in np.arange(0, self.num_neurons, self.ONE_TIME_FILTERS): Nfilters = min(self.ONE_TIME_FILTERS, self.num_neurons - i) self.generate_filters(startbias=i, N_filters=Nfilters) la.dot(self.filters, d_image, opb='t', C=d_output[i: i+Nfilters], handle=handle) del self.filters return d_output.T()
def __init__(self, decode_neuron, tk1, tk2,neuron_ind, kappa, delta, bias, norm, Wt, Mt, dtype = np.float64): """ class dirichlet For use to reconstruct videos encoded by VTEMs consist of spatial only receptive fields and IAF neurons parameters: decode_neuron: number of neuron used in the decoding tk1: a ndarray vector of float64 containing all the t_k's in \int_{t_k}^{t_{k+1}} tk2: a ndarray vector of float64 containing all the t_{k+1}'s in \int_{t_k}^{k+1}} neuron_ind: a ndarray vector of int32, where neuron_ind[j] is the neuron index correspond to (tk1[j],tk2[j]) kappa: a ndarray vector of float64 containing the integration constants for each neuron delta: a ndarray vector of float64 containing the threshold for each neuron bias: a ndarray vector of float64 containing the bias for each neuron norm: a ndarray vector of float64 containing the normalization factor for each neuron (in case of deterministic threshold, use a vector of ones) Wt: temporal bandwidth Mt: dirichlet order in variable t dtype: np.float64 or np.float32, determine dtype of G matrix and q vector """ self.size = tk1.size if self.size == 0: raise ValueError("no spikes to decode") self.Wt = Wt self.Mt = Mt self.dtype = np.dtype(dtype) print "number of spikes: %d" % (self.size) self.d_tk1 = parray.to_gpu(tk1) self.d_tk2 = parray.to_gpu(tk2) self.d_neuron_ind = parray.to_gpu(neuron_ind) self.d_kappa = parray.to_gpu(kappa) self.d_delta = parray.to_gpu(delta) self.d_bias = parray.to_gpu(bias) self.d_norm = parray.to_gpu(norm)
def _init_objects(self): self.neurons = [ self._instantiate_neuron(i, t, n) for i, (t, n) in enumerate(self.n_list) if t!=PORT_IN_GPOT and t!=PORT_IN_SPK] self.synapses = [ self._instantiate_synapse(i, t, n) for i, (t, n) in enumerate(self.s_list) if t!='pass'] self.buffer = CircularArray(self.total_num_gpot_neurons, self.gpot_delay_steps, self.V, self.total_num_spike_neurons, self.spike_delay_steps) if self.input_file: self.input_h5file = tables.openFile(self.input_file) self.file_pointer = 0 self.I_ext = parray.to_gpu(self.input_h5file.root.array.read( self.file_pointer, self.file_pointer + self._one_time_import)) self.file_pointer += self._one_time_import self.frame_count = 0 self.frames_in_buffer = self._one_time_import if self.output: output_file = self.output_file.rsplit('.', 1) filename = output_file[0] if len(output_file) > 1: ext = output_file[1] else: ext = 'h5' if self.total_num_gpot_neurons > 0: self.output_gpot_file = tables.openFile(filename + '_gpot.' + ext, 'w') self.output_gpot_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.total_num_gpot_neurons)) if self.total_num_spike_neurons > 0: self.output_spike_file = tables.openFile(filename + '_spike.' + ext, 'w') self.output_spike_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.total_num_spike_neurons)) if self.debug: ''' self.in_gpot_files = {} for (key, i) in self.other_lpu_map.iteritems(): num = self.num_input_gpot_neurons[i] if num>0: self.in_gpot_files[key] = tables.openFile(filename + \ key + '_in_gpot.' + ext , 'w') self.in_gpot_files[key].createEArray("/", "array", \ tables.Float64Atom(), (0, num)) ''' if self.total_num_gpot_neurons > 0: self.gpot_buffer_file = tables.openFile(self.id + '_buffer.h5','w') self.gpot_buffer_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.gpot_delay_steps, self.total_num_gpot_neurons)) if self.total_synapses + len(self.input_neuron_list) > 0: self.synapse_state_file = tables.openFile(self.id + '_synapses.h5', 'w') self.synapse_state_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.total_synapses + len(self.input_neuron_list)))
def _init_objects(self): self.neurons = [ self._instantiate_neuron(i, t, n) for i, (t, n) in enumerate(self.n_list) if t!=PORT_IN_GPOT and t!=PORT_IN_SPK] self.synapses = [ self._instantiate_synapse(i, t, n) for i, (t, n) in enumerate(self.s_list) if t!='pass'] self.buffer = CircularArray(self.total_num_gpot_neurons, self.gpot_delay_steps, self.V, self.total_num_spike_neurons, self.spike_delay_steps) if self.input_file: self.input_h5file = h5py.File(self.input_file, 'r') self.file_pointer = 0 self.I_ext = \ parray.to_gpu(self.input_h5file['/array'][self.file_pointer: self.file_pointer+self._one_time_import]) self.file_pointer += self._one_time_import self.frame_count = 0 self.frames_in_buffer = self._one_time_import if self.output: output_file = self.output_file.rsplit('.', 1) filename = output_file[0] if len(output_file) > 1: ext = output_file[1] else: ext = 'h5' if self.total_num_gpot_neurons > 0: self.output_gpot_file = h5py.File(filename+'_gpot.'+ext, 'w') self.output_gpot_file.create_dataset( '/array', (0, self.total_num_gpot_neurons), dtype=np.float64, maxshape=(None, self.total_num_gpot_neurons)) if self.total_num_spike_neurons > 0: self.output_spike_file = h5py.File(filename+'_spike.'+ext, 'w') self.output_spike_file.create_dataset( '/array', (0, self.total_num_spike_neurons), dtype=np.float64, maxshape=(None, self.total_num_spike_neurons)) if self.debug: if self.total_num_gpot_neurons > 0: self.gpot_buffer_file = h5py.File(self.id + '_buffer.h5', 'w') self.gpot_buffer_file.create_dataset( '/array', (0, self.gpot_delay_steps, self.total_num_gpot_neurons), dtype=np.float64, maxshape=(None, self.gpot_delay_steps, self.total_num_gpot_neurons)) if self.total_synapses + len(self.input_neuron_list) > 0: self.synapse_state_file = h5py.File(self.id + '_synapses.h5', 'w') self.synapse_state_file.create_dataset( '/array', (0, self.total_synapses + len(self.input_neuron_list)), dtype=np.float64, maxshape=(None, self.total_synapses + len(self.input_neuron_list)))
def restore_others(self): """ restore GPU memory with tk1, tk2, neuron_ind, norm """ self.d_tk1 = parray.to_gpu(self.tk1) self.d_tk2 = parray.to_gpu(self.tk2) self.d_neuron_ind = parray.to_gpu(self.neuron_ind) self.d_norm = parray.to_gpu(self.norm)
def restore_Gq(self): """ restore GPU memory with G, q """ self.d_G = parray.to_gpu(self.G) self.d_q = parray.to_gpu(self.q)
def to_gpu(self): self.d_refz = parray.to_gpu(self.refz) self.d_reftheta = parray.to_gpu(self.reftheta) self.d_grid = [parray.to_gpu(self.grid[i].reshape(-1)) for i in range(len(self.grid))]
def to_gpu(self): self.d_refelev = parray.to_gpu(self.refelev) self.d_refazim = parray.to_gpu(self.refazim) self.d_grid = [parray.to_gpu(self.grid[i].reshape(-1)) for i in range(len(self.grid))]
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
def _init_objects(self): self.neurons = [ self._instantiate_neuron(i, t, n) for i, (t, n) in enumerate(self.n_list) if t != PORT_IN_GPOT and t != PORT_IN_SPK ] self.synapses = [self._instantiate_synapse(i, t, n) for i, (t, n) in enumerate(self.s_list) if t != "pass"] self.buffer = CircularArray( self.total_num_gpot_neurons, self.gpot_delay_steps, self.V, self.total_num_spike_neurons, self.spike_delay_steps, ) if self.input_file: self.input_h5file = tables.openFile(self.input_file) self.file_pointer = 0 self.I_ext = parray.to_gpu( self.input_h5file.root.array.read(self.file_pointer, self.file_pointer + self._one_time_import) ) self.file_pointer += self._one_time_import self.frame_count = 0 self.frames_in_buffer = self._one_time_import if self.output: output_file = self.output_file.rsplit(".", 1) filename = output_file[0] if len(output_file) > 1: ext = output_file[1] else: ext = "h5" if self.total_num_gpot_neurons > 0: self.output_gpot_file = tables.openFile(filename + "_gpot." + ext, "w") self.output_gpot_file.createEArray("/", "array", tables.Float64Atom(), (0, self.total_num_gpot_neurons)) if self.total_num_spike_neurons > 0: self.output_spike_file = tables.openFile(filename + "_spike." + ext, "w") self.output_spike_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.total_num_spike_neurons) ) if self.debug: """ self.in_gpot_files = {} for (key, i) in self.other_lpu_map.iteritems(): num = self.num_input_gpot_neurons[i] if num>0: self.in_gpot_files[key] = tables.openFile(filename + \ key + '_in_gpot.' + ext , 'w') self.in_gpot_files[key].createEArray("/", "array", \ tables.Float64Atom(), (0, num)) """ if self.total_num_gpot_neurons > 0: self.gpot_buffer_file = tables.openFile(self.id + "_buffer.h5", "w") self.gpot_buffer_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.gpot_delay_steps, self.total_num_gpot_neurons) ) if self.total_synapses + len(self.input_neuron_list) > 0: self.synapse_state_file = tables.openFile(self.id + "_synapses.h5", "w") self.synapse_state_file.createEArray( "/", "array", tables.Float64Atom(), (0, self.total_synapses + len(self.input_neuron_list)) )
S1 = 128 S2 = 128 PHOTORECEPTORS = 8 M_size = S1*S2 # same as grid[0].size N_filters = PHOTORECEPTORS RAD = 1 KAPPA = 20 SIGMA = 1 # or angle NTHREADS = (128, 1, 1) NBLOCKS = ((M_size-1) // NTHREADS[0] + 1, 1) d_filters = parray.empty((N_filters, M_size), dtype) grid = np.meshgrid(np.linspace(-1, 1, num=S1), np.linspace(-np.pi, np.pi, num=S2)) d_grid = [parray.to_gpu(grid[i].flatten()) for i in range(len(grid))] dxy = np.diff(grid[0][0, :2])*np.diff(grid[1][:2, 0])[0] ref_z = 2*random.rand(PHOTORECEPTORS)-1 # -1 to 1 d_refz = parray.to_gpu(ref_z) ref_theta = np.pi*random.rand(PHOTORECEPTORS) # half cylinder d_reftheta = parray.to_gpu(ref_theta) filter_func.prepared_call( NBLOCKS, NTHREADS, d_filters.gpudata, d_filters.ld, d_grid[0].gpudata, d_grid[1].gpudata,