def test_hermitian_complex128(self):
# M < N
a = np.array([[1j, 2j, 3j, 4j, 5j, 6j], [7j, 8j, 9j, 10j, 11j, 12j]], np.complex128)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(np.conj(a.T) == at_gpu.get())
# M > N
b = a.T.copy()
b_gpu = gpuarray.to_gpu(b)
bt_gpu = linalg.hermitian(b_gpu)
assert np.all(np.conj(b.T) == bt_gpu.get())
def test_hermitian_float64(self):
# M < N
a = np.array([[1, 2, 3, 4, 5, 6], [7, 8, 9, 10, 11, 12]], np.float64)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(a.T == at_gpu.get())
# M > N
b = a.T.copy()
b_gpu = gpuarray.to_gpu(b)
bt_gpu = linalg.hermitian(b_gpu)
assert np.all(b.T == bt_gpu.get())
def test_hermitian_float64(self):
# M < N
a = np.array([[1, 2, 3, 4, 5, 6], [7, 8, 9, 10, 11, 12]], np.float64)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(a.T == at_gpu.get())
# M > N
b = a.T.copy()
b_gpu = gpuarray.to_gpu(b)
bt_gpu = linalg.hermitian(b_gpu)
assert np.all(b.T == bt_gpu.get())
def test_hermitian_complex128(self):
# M < N
a = np.array([[1j, 2j, 3j, 4j, 5j, 6j], [7j, 8j, 9j, 10j, 11j, 12j]],
np.complex128)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(np.conj(a.T) == at_gpu.get())
# M > N
b = a.T.copy()
b_gpu = gpuarray.to_gpu(b)
bt_gpu = linalg.hermitian(b_gpu)
assert np.all(np.conj(b.T) == bt_gpu.get())
def test_hermitian_float64(self):
a = np.array([[1, 2, 3, 4, 5, 6],
[7, 8, 9, 10, 11, 12]],
np.float64)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(a.T == at_gpu.get())
def test_hermitian_complex64(self):
a = np.array([[1j, 2j, 3j, 4j, 5j, 6j],
[7j, 8j, 9j, 10j, 11j, 12j]],
np.complex64)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(np.conj(a.T) == at_gpu.get())
def iaf_decode_pop(s_gpu, ns_gpu, dur, dt, bw, b_gpu, d_gpu, R_gpu,
C_gpu, M=5, smoothing=0.0):
"""
Population IAF time decoding machine.
Decode a signal encoded with an ensemble of Integrate-and-Fire
neurons assuming that the encoded signal is representable in terms
of trigonometric polynomials.
Parameters
----------
s_gpu : pycuda.gpuarray.GPUArray
Signal encoded by an ensemble of encoders. The nonzero
values represent the time between spikes (in s). The number of
arrays in the list corresponds to the number of encoders in
the ensemble.
ns_gpu : pycuda.gpuarray.GPUArray
Number of interspike intervals in each row of `s_gpu`.
dur : float
Duration of signal (in s).
dt : float
Sampling resolution of original signal; the sampling frequency
is 1/dt Hz.
bw : float
Signal bandwidth (in rad/s).
b_gpu : pycuda.gpuarray.GPUArray
Array of encoder biases.
d_gpu : pycuda.gpuarray.GPUArray
Array of encoder thresholds.
R_gpu : pycuda.gpuarray.GPUArray
Array of neuron resistances.
C_gpu : pycuda.gpuarray.GPUArray
Array of neuron capacitances.
M : int
2*M+1 coefficients are used for reconstructing the signal.
smoothing : float
Smoothing parameter.
Returns
-------
u_rec : pycuda.gpuarray.GPUArray
Recovered signal.
Notes
-----
The number of spikes contributed by each neuron may differ from the
number contributed by other neurons.
"""
# Sanity checks:
float_type = s_gpu.dtype.type
if float_type == np.float32:
use_double = 0
complex_type = np.complex64
__pinv_rcond__ = 1e-4
elif float_type == np.float64:
use_double = 1
complex_type = np.complex128
__pinv_rcond__ = 1e-8
else:
raise ValueError('unsupported data type')
N = s_gpu.shape[0]
if not N:
raise ValueError('no spike data given')
if (ns_gpu.size != N) or (b_gpu.size != N) or (d_gpu.size != N) or \
(R_gpu.size != N) or (C_gpu.size != N):
raise ValueError('parameter arrays must be of same length')
T = 2*np.pi*M/bw
if T < dur:
raise ValueError('2*pi*M/bw must exceed the signal length')
# Map CUDA index to neuron index and interspike interval index:
ns = ns_gpu.get()
idx_to_ni, idx_to_k = _compute_idx_map(ns)
idx_to_ni_gpu = gpuarray.to_gpu(idx_to_ni)
idx_to_k_gpu = gpuarray.to_gpu(idx_to_k)
dev = cumisc.get_current_device()
# Use a smaller block size than the maximum to prevent the kernels
# from using too many registers:
max_threads_per_block = 256
# Prepare kernels:
cache_dir = None
compute_ts_pop_mod = SourceModule(compute_ts_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir)
compute_ts_pop = compute_ts_pop_mod.get_function('compute_ts')
compute_q_pop_mod = \
SourceModule(compute_q_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir)
compute_q_pop_ideal = compute_q_pop_mod.get_function('compute_q_ideal')
compute_q_pop_leaky = compute_q_pop_mod.get_function('compute_q_leaky')
compute_F_pop_mod = \
SourceModule(compute_F_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir,
options=['-I', install_headers])
compute_F_pop_ideal = compute_F_pop_mod.get_function('compute_F_ideal')
compute_F_pop_leaky = compute_F_pop_mod.get_function('compute_F_leaky')
compute_u_pop_mod = \
SourceModule(compute_u_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir,
options=['-I', install_headers])
compute_u_pop = compute_u_pop_mod.get_function('compute_u')
# Total number of interspike intervals per neuron less 1 for each
# spike train with more than
Nq = int(np.sum(ns)-np.sum(ns>1))
# Set up GPUArrays for intermediary data:
ts_gpu = gpuarray.zeros_like(s_gpu)
# Note that these arrays are complex to enable use of CUBLAS
# matrix multiplication functions:
q_gpu = gpuarray.empty((Nq, 1), complex_type)
F_gpu = gpuarray.empty((Nq, 2*M+1), complex_type)
# Get required block/grid sizes:
block_dim_ts, grid_dim_ts = \
cumisc.select_block_grid_sizes(dev, N,
max_threads_per_block)
block_dim_q, grid_dim_q = \
cumisc.select_block_grid_sizes(dev, q_gpu.shape,
max_threads_per_block)
block_dim_F, grid_dim_F = \
cumisc.select_block_grid_sizes(dev, F_gpu.shape,
max_threads_per_block)
# Launch kernels:
compute_ts_pop(s_gpu, ns_gpu, ts_gpu, np.uint32(s_gpu.shape[1]),
np.uint32(N),
block=block_dim_ts, grid=grid_dim_ts)
if np.all(np.isinf(R_gpu.get())):
compute_q_pop_ideal(s_gpu, q_gpu,
b_gpu, d_gpu, C_gpu,
idx_to_ni_gpu, idx_to_k_gpu,
np.uint32(s_gpu.shape[1]),
np.uint32(Nq),
block=block_dim_q, grid=grid_dim_q)
compute_F_pop_ideal(s_gpu, ts_gpu, F_gpu,
float_type(bw),
idx_to_ni_gpu, idx_to_k_gpu,
np.int32(M), np.uint32(s_gpu.shape[1]),
np.uint32(F_gpu.size),
block=block_dim_F, grid=grid_dim_F)
else:
compute_q_pop_leaky(s_gpu, q_gpu,
b_gpu, d_gpu,
R_gpu, C_gpu,
idx_to_ni_gpu, idx_to_k_gpu,
np.uint32(s_gpu.shape[1]),
np.uint32(Nq),
block=block_dim_q, grid=grid_dim_q)
compute_F_pop_leaky(s_gpu, ts_gpu, F_gpu,
float_type(bw), R_gpu, C_gpu,
idx_to_ni_gpu, idx_to_k_gpu,
np.int32(M), np.uint32(s_gpu.shape[1]),
np.uint32(F_gpu.size),
block=block_dim_F, grid=grid_dim_F)
# Free unneeded variables:
del s_gpu, ts_gpu, idx_to_ni_gpu, idx_to_k_gpu
# Compute the product of F^H and q first so that both F^H and q
# can be dropped from memory:
FH_gpu = culinalg.hermitian(F_gpu)
FHq_gpu = culinalg.dot(FH_gpu, q_gpu)
del FH_gpu, q_gpu
if smoothing == 0:
c_gpu = culinalg.dot(culinalg.pinv(culinalg.dot(F_gpu, F_gpu, 'c'),
__pinv_rcond__),
FHq_gpu)
else:
c_gpu = culinalg.dot(culinalg.pinv(culinalg.dot(F_gpu, F_gpu, 'c')+
np.sum(ns)*smoothing*culinalg.eye(2*M+1,
float_type),
__pinv_rcond__),
FHq_gpu)
# Allocate array for reconstructed signal:
Nt = int(np.ceil(dur/dt))
u_rec_gpu = gpuarray.to_gpu(np.zeros(Nt, complex_type))
### Replace the above with the following line when the bug in
# gpuarray.zeros in pycuda 2011.1.2 is fixed:
#u_rec_gpu = gpuarray.zeros(Nt, complex_type)
# Get required block/grid sizes:
block_dim_t, grid_dim_t = \
cumisc.select_block_grid_sizes(dev, Nt, max_threads_per_block)
# Reconstruct signal:
compute_u_pop(u_rec_gpu, c_gpu, float_type(bw),
float_type(dt),
np.int32(M),
np.uint32(Nt),
block=block_dim_t, grid=grid_dim_t)
return np.real(u_rec_gpu.get())
def test_hermitian_float64(self):
a = np.array([[1, 2, 3, 4, 5, 6], [7, 8, 9, 10, 11, 12]], np.float64)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(a.T == at_gpu.get())
def test_hermitian_complex64(self):
a = np.array([[1j, 2j, 3j, 4j, 5j, 6j], [7j, 8j, 9j, 10j, 11j, 12j]],
np.complex64)
a_gpu = gpuarray.to_gpu(a)
at_gpu = linalg.hermitian(a_gpu)
assert np.all(np.conj(a.T) == at_gpu.get())
def iaf_decode_pop(s_gpu,
ns_gpu,
dur,
dt,
bw,
b_gpu,
d_gpu,
R_gpu,
C_gpu,
M=5,
smoothing=0.0):
"""
Population IAF time decoding machine.
Decode a signal encoded with an ensemble of Integrate-and-Fire
neurons assuming that the encoded signal is representable in terms
of trigonometric polynomials.
Parameters
----------
s_gpu : pycuda.gpuarray.GPUArray
Signal encoded by an ensemble of encoders. The nonzero
values represent the time between spikes (in s). The number of
arrays in the list corresponds to the number of encoders in
the ensemble.
ns_gpu : pycuda.gpuarray.GPUArray
Number of interspike intervals in each row of `s_gpu`.
dur : float
Duration of signal (in s).
dt : float
Sampling resolution of original signal; the sampling frequency
is 1/dt Hz.
bw : float
Signal bandwidth (in rad/s).
b_gpu : pycuda.gpuarray.GPUArray
Array of encoder biases.
d_gpu : pycuda.gpuarray.GPUArray
Array of encoder thresholds.
R_gpu : pycuda.gpuarray.GPUArray
Array of neuron resistances.
C_gpu : pycuda.gpuarray.GPUArray
Array of neuron capacitances.
M : int
2*M+1 coefficients are used for reconstructing the signal.
smoothing : float
Smoothing parameter.
Returns
-------
u_rec : pycuda.gpuarray.GPUArray
Recovered signal.
Notes
-----
The number of spikes contributed by each neuron may differ from the
number contributed by other neurons.
"""
# Sanity checks:
float_type = s_gpu.dtype.type
if float_type == np.float32:
use_double = 0
complex_type = np.complex64
__pinv_rcond__ = 1e-4
elif float_type == np.float64:
use_double = 1
complex_type = np.complex128
__pinv_rcond__ = 1e-8
else:
raise ValueError('unsupported data type')
N = s_gpu.shape[0]
if not N:
raise ValueError('no spike data given')
if (ns_gpu.size != N) or (b_gpu.size != N) or (d_gpu.size != N) or \
(R_gpu.size != N) or (C_gpu.size != N):
raise ValueError('parameter arrays must be of same length')
T = 2 * np.pi * M / bw
if T < dur:
raise ValueError('2*pi*M/bw must exceed the signal length')
# Map CUDA index to neuron index and interspike interval index:
ns = ns_gpu.get()
idx_to_ni, idx_to_k = _compute_idx_map(ns)
idx_to_ni_gpu = gpuarray.to_gpu(idx_to_ni)
idx_to_k_gpu = gpuarray.to_gpu(idx_to_k)
dev = cumisc.get_current_device()
# Use a smaller block size than the maximum to prevent the kernels
# from using too many registers:
max_threads_per_block = 256
# Prepare kernels:
cache_dir = None
compute_ts_pop_mod = SourceModule(
compute_ts_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir)
compute_ts_pop = compute_ts_pop_mod.get_function('compute_ts')
compute_q_pop_mod = \
SourceModule(compute_q_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir)
compute_q_pop_ideal = compute_q_pop_mod.get_function('compute_q_ideal')
compute_q_pop_leaky = compute_q_pop_mod.get_function('compute_q_leaky')
compute_F_pop_mod = \
SourceModule(compute_F_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir,
options=['-I', install_headers])
compute_F_pop_ideal = compute_F_pop_mod.get_function('compute_F_ideal')
compute_F_pop_leaky = compute_F_pop_mod.get_function('compute_F_leaky')
compute_u_pop_mod = \
SourceModule(compute_u_pop_template.substitute(use_double=use_double),
cache_dir=cache_dir,
options=['-I', install_headers])
compute_u_pop = compute_u_pop_mod.get_function('compute_u')
# Total number of interspike intervals per neuron less 1 for each
# spike train with more than
Nq = int(np.sum(ns) - np.sum(ns > 1))
# Set up GPUArrays for intermediary data:
ts_gpu = gpuarray.zeros_like(s_gpu)
# Note that these arrays are complex to enable use of CUBLAS
# matrix multiplication functions:
q_gpu = gpuarray.empty((Nq, 1), complex_type)
F_gpu = gpuarray.empty((Nq, 2 * M + 1), complex_type)
# Get required block/grid sizes:
block_dim_ts, grid_dim_ts = \
cumisc.select_block_grid_sizes(dev, N,
max_threads_per_block)
block_dim_q, grid_dim_q = \
cumisc.select_block_grid_sizes(dev, q_gpu.shape,
max_threads_per_block)
block_dim_F, grid_dim_F = \
cumisc.select_block_grid_sizes(dev, F_gpu.shape,
max_threads_per_block)
# Launch kernels:
compute_ts_pop(s_gpu,
ns_gpu,
ts_gpu,
np.uint32(s_gpu.shape[1]),
np.uint32(N),
block=block_dim_ts,
grid=grid_dim_ts)
if np.all(np.isinf(R_gpu.get())):
compute_q_pop_ideal(s_gpu,
q_gpu,
b_gpu,
d_gpu,
C_gpu,
idx_to_ni_gpu,
idx_to_k_gpu,
np.uint32(s_gpu.shape[1]),
np.uint32(Nq),
block=block_dim_q,
grid=grid_dim_q)
compute_F_pop_ideal(s_gpu,
ts_gpu,
F_gpu,
float_type(bw),
idx_to_ni_gpu,
idx_to_k_gpu,
np.int32(M),
np.uint32(s_gpu.shape[1]),
np.uint32(F_gpu.size),
block=block_dim_F,
grid=grid_dim_F)
else:
compute_q_pop_leaky(s_gpu,
q_gpu,
b_gpu,
d_gpu,
R_gpu,
C_gpu,
idx_to_ni_gpu,
idx_to_k_gpu,
np.uint32(s_gpu.shape[1]),
np.uint32(Nq),
block=block_dim_q,
grid=grid_dim_q)
compute_F_pop_leaky(s_gpu,
ts_gpu,
F_gpu,
float_type(bw),
R_gpu,
C_gpu,
idx_to_ni_gpu,
idx_to_k_gpu,
np.int32(M),
np.uint32(s_gpu.shape[1]),
np.uint32(F_gpu.size),
block=block_dim_F,
grid=grid_dim_F)
# Free unneeded variables:
del s_gpu, ts_gpu, idx_to_ni_gpu, idx_to_k_gpu
# Compute the product of F^H and q first so that both F^H and q
# can be dropped from memory:
FH_gpu = culinalg.hermitian(F_gpu)
FHq_gpu = culinalg.dot(FH_gpu, q_gpu)
del FH_gpu, q_gpu
if smoothing == 0:
c_gpu = culinalg.dot(
culinalg.pinv(culinalg.dot(F_gpu, F_gpu, 'c'), __pinv_rcond__),
FHq_gpu)
else:
c_gpu = culinalg.dot(
culinalg.pinv(
culinalg.dot(F_gpu, F_gpu, 'c') +
np.sum(ns) * smoothing * culinalg.eye(2 * M + 1, float_type),
__pinv_rcond__), FHq_gpu)
# Allocate array for reconstructed signal:
Nt = int(np.ceil(dur / dt))
u_rec_gpu = gpuarray.to_gpu(np.zeros(Nt, complex_type))
### Replace the above with the following line when the bug in
# gpuarray.zeros in pycuda 2011.1.2 is fixed:
#u_rec_gpu = gpuarray.zeros(Nt, complex_type)
# Get required block/grid sizes:
block_dim_t, grid_dim_t = \
cumisc.select_block_grid_sizes(dev, Nt, max_threads_per_block)
# Reconstruct signal:
compute_u_pop(u_rec_gpu,
c_gpu,
float_type(bw),
float_type(dt),
np.int32(M),
np.uint32(Nt),
block=block_dim_t,
grid=grid_dim_t)
return np.real(u_rec_gpu.get())