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 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 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 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,
