Skip to content

Commit

Permalink
added support to run on 32-bit machines
Browse files Browse the repository at this point in the history
  • Loading branch information
@yiyin
yiyin committed Dec 24, 2015
1 parent 82c8fdf commit a4c80c8
Show file tree
Hide file tree
Showing 18 changed files with 2,937 additions and 1,089 deletions.
2 changes: 1 addition & 1 deletion neurokernel/LPU/LPU.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -970,7 +970,7 @@ def _extract_projection_func(self, state_var):
template % {"type": dtype_to_ctype(state_var.dtype)},
options=self.compile_options)
func = mod.get_function("extract_projection")
func.prepare([np.intp, np.intp, np.intp, np.intp, np.int32])
func.prepare('PPPPi')#[np.intp, np.intp, np.intp, np.intp, np.int32])

return func

Expand Down
24 changes: 13 additions & 11 deletions neurokernel/LPU/neurons/LeakyIAF.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -131,15 +131,16 @@ def get_gpu_kernel( self):
"nneu": self.gpu_block[0] },
options=["--ptxas-options=-v"])
func = mod.get_function("leaky_iaf")
func.prepare( [ np.int32, # neu_num
np.float64, # dt
np.intp, # spk array
np.intp, # V array
np.intp, # I array
np.intp, # Vt array
np.intp, # Vr array
np.intp, # R array
np.intp ]) # C array
func.prepare('idPPPPPPP')
# [ np.int32, # neu_num
# np.float64, # dt
# np.intp, # spk array
# np.intp, # V array
# np.intp, # I array
# np.intp, # Vt array
# np.intp, # Vr array
# np.intp, # R array
# np.intp ]) # C array

return func

Expand Down Expand Up @@ -266,7 +267,8 @@ def _get_update_I_cond_func(self):
mod = SourceModule(template % {"num_neurons": self.num_neurons},
options = ["--ptxas-options=-v"])
func = mod.get_function("get_input")
func.prepare([np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PPPPPPP')
#[np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
self._block_get_input = (32,32,1)
self._grid_get_input = ((self.num_neurons - 1) / 32 + 1, 1)
return func
Expand Down Expand Up @@ -361,5 +363,5 @@ def _get_update_I_non_cond_func(self):
mod = SourceModule(template % {"num_neurons": self.num_neurons},
options = ["--ptxas-options=-v"])
func = mod.get_function("get_input")
func.prepare([np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PPPPP')#[np.intp, np.intp, np.intp, np.intp, np.intp])
return func
26 changes: 14 additions & 12 deletions neurokernel/LPU/neurons/LeakyIAF_bias.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -125,16 +125,17 @@ def get_gpu_kernel( self):
"nneu": self.gpu_block[0] },\
options=["--ptxas-options=-v"])
func = mod.get_function("leaky_iaf")
func.prepare( [ np.int32, # neu_num
np.float64, # dt
np.intp, # spk array
np.intp, # V array
np.intp, # I array
np.intp, # Vt array
np.intp, # Vr array
np.intp, # R array
np.intp, # C array
np.intp]) # b array
func.prepare('idPPPPPPPP')
# [ np.int32, # neu_num
# np.float64, # dt
# np.intp, # spk array
# np.intp, # V array
# np.intp, # I array
# np.intp, # Vt array
# np.intp, # Vr array
# np.intp, # R array
# np.intp, # C array
# np.intp]) # b array

return func

Expand Down Expand Up @@ -251,7 +252,8 @@ def _get_update_I_cond_func(self):
"""
mod = SourceModule(template % {"num_neurons": self.num_neurons}, options = ["--ptxas-options=-v"])
func = mod.get_function("get_input")
func.prepare([np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PPPPPPP')
#[np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
self._block_get_input = (32,32,1)
self._grid_get_input = ((self.num_neurons - 1) / 32 + 1, 1)
return func
Expand Down Expand Up @@ -344,5 +346,5 @@ def _get_update_I_non_cond_func(self):
"""
mod = SourceModule(template % {"num_neurons": self.num_neurons}, options = ["--ptxas-options=-v"])
func = mod.get_function("get_input")
func.prepare([np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PPPPP')#[np.intp, np.intp, np.intp, np.intp, np.intp])
return func
7 changes: 4 additions & 3 deletions neurokernel/LPU/neurons/MorrisLecar.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -127,9 +127,10 @@ def get_euler_kernel(self):
func = mod.get_function("hhn_euler_multiple")


func.prepare([np.intp, np.intp, np.int32, np.intp, scalartype,
np.int32, np.intp, np.intp, np.intp,
np.intp, np.intp, np.intp])
func.prepare('PPiP'+np.dtype(scalartype).char+'iPPPPPP')
#[np.intp, np.intp, np.int32, np.intp, scalartype,
# np.int32, np.intp, np.intp, np.intp,
# np.intp, np.intp, np.intp])


return func
4 changes: 3 additions & 1 deletion neurokernel/LPU/neurons/MorrisLecarCopy.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -116,7 +116,9 @@ def get_euler_kernel(self):
func = mod.get_function("hhn_euler_multiple")


func.prepare([np.intp, np.intp, np.int32, np.intp, scalartype, np.int32, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PPiP'+np.dtype(scalartype).char+'iPPPPPP')
#[np.intp, np.intp, np.int32, np.intp, scalartype, np.int32, np.intp, np.intp,
# np.intp, np.intp, np.intp, np.intp])


return func
10 changes: 5 additions & 5 deletions neurokernel/LPU/neurons/MorrisLecar_a.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -126,10 +126,10 @@ def get_euler_kernel(self):
func = mod.get_function("hhn_euler_multiple")


func.prepare([np.intp, np.intp, np.int32, np.intp, scalartype,
np.int32, np.intp, np.intp, np.intp, np.intp,
np.intp, np.intp, np.intp, np.intp, np.intp,
np.intp, np.intp, np.intp])

func.prepare('PPiP'+np.dtype(scalartype).char+'iPPPPPPPPPPPP')
# [np.intp, np.intp, np.int32, np.intp, scalartype,
# np.int32, np.intp, np.intp, np.intp, np.intp,
# np.intp, np.intp, np.intp, np.intp, np.intp,
# np.intp, np.intp, np.intp])

return func
23 changes: 12 additions & 11 deletions neurokernel/LPU/synapses/AlphaSynapse.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -101,15 +101,16 @@ def get_gpu_kernel(self):
cuda_src % {"type": dtype_to_ctype(np.float64)},\
options=["--ptxas-options=-v"])
func = mod.get_function("alpha_synapse")
func.prepare( [ np.int32, # syn_num
np.float64, # dt
np.intp, # spike list
np.intp, # pre-synaptic neuron list
np.intp, # ar array
np.intp, # ad array
np.intp, # gmax array
np.intp, # a0 array
np.intp, # a1 array
np.intp, # a2 array
np.intp ] ) # cond array
func.prepare('idPPPPPPPPP')
# [ np.int32, # syn_num
# np.float64, # dt
# np.intp, # spike list
# np.intp, # pre-synaptic neuron list
# np.intp, # ar array
# np.intp, # ad array
# np.intp, # gmax array
# np.intp, # a0 array
# np.intp, # a1 array
# np.intp, # a2 array
# np.intp ] ) # cond array
return func
36 changes: 19 additions & 17 deletions neurokernel/LPU/synapses/AlphaSynapsePre.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -230,30 +230,32 @@ def _get_gpu_kernel(self):
cuda_src_synapse_kernel % {"type": dtype_to_ctype(np.float64)},\
options=["--ptxas-options=-v"])
func = mod.get_function("alpha_synapse")
func.prepare([np.int32, # syn_num
np.float64, # dt
np.intp, # spike list
np.intp, # pre-synaptic neuron list
np.intp, # ar array
np.intp, # ad array
np.intp, # gmax array
np.intp, # a0 array
np.intp, # a1 array
np.intp, # a2 array
np.intp, # pre-synaptic input
np.intp]) # cond array
func.prepare('idPPPPPPPPPP')
# [np.int32, # syn_num
# np.float64, # dt
# np.intp, # spike list
# np.intp, # pre-synaptic neuron list
# np.intp, # ar array
# np.intp, # ad array
# np.intp, # gmax array
# np.intp, # a0 array
# np.intp, # a1 array
# np.intp, # a2 array
# np.intp, # pre-synaptic input
# np.intp]) # cond array
return func

def _get_update_I_non_cond_func(self):
mod = SourceModule(\
cuda_src_synapse_update_I % {"num": self.num},
options = ["--ptxas-options=-v"])
func = mod.get_function("get_input")
func.prepare([np.intp, # synapse state
np.intp, # cumulative dendrites number
np.intp, # dendrites number
np.intp, # pre-synaptic number ID
np.intp]) # output
func.prepare('PPPPP')
# [np.intp, # synapse state
# np.intp, # cumulative dendrites number
# np.intp, # dendrites number
# np.intp, # pre-synaptic number ID
# np.intp]) # output

self._block_get_input = (32,32,1)
self._grid_get_input = ((self.num - 1) / 32 + 1, 1)
Expand Down
17 changes: 9 additions & 8 deletions neurokernel/LPU/synapses/DummySynapse.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -77,12 +77,13 @@ def get_gpu_kernel(self):
cuda_src % {"type": dtype_to_ctype(np.float64)},\
options=["--ptxas-options=-v"])
func = mod.get_function("dummy_synapse")
func.prepare( [ np.intp, # neuron state buffer
np.int32, # buffer width
np.int32, # buffer position
np.int32, # buffer delay steps
np.int32, # syn_num
np.intp, # pre-synaptic neuron list
np.intp, # delay step
np.intp ] ) # cond array
func.prepare('PiiiiPPP')
# [ np.intp, # neuron state buffer
# np.int32, # buffer width
# np.int32, # buffer position
# np.int32, # buffer delay steps
# np.int32, # syn_num
# np.intp, # pre-synaptic neuron list
# np.intp, # delay step
# np.intp ] ) # cond array
return func
19 changes: 10 additions & 9 deletions neurokernel/LPU/synapses/ExpSynapse.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -93,13 +93,14 @@ def _get_gpu_kernel(self):
cuda_src % {"type": dtype_to_ctype(np.float64)},\
options=["--ptxas-options=-v"])
func = mod.get_function("exponential_synapse")
func.prepare( [ np.int32, # syn_num
np.float64, # dt
np.intp, # spike list
np.intp, # pre-synaptic neuron list
np.intp, # tau; time constant
np.intp, # a; bump size
np.intp, # gmax array
np.intp, # eff; efficacy
np.intp ] ) # cond array
func.prepare('idPPPPPPP')
# [ np.int32, # syn_num
# np.float64, # dt
# np.intp, # spike list
# np.intp, # pre-synaptic neuron list
# np.intp, # tau; time constant
# np.intp, # a; bump size
# np.intp, # gmax array
# np.intp, # eff; efficacy
# np.intp ] ) # cond array
return func
30 changes: 16 additions & 14 deletions neurokernel/LPU/synapses/ExpSynapsePre.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -214,27 +214,29 @@ def _get_gpu_kernel(self):
cuda_src_synapse_kernel % {"type": dtype_to_ctype(np.float64)},\
options=["--ptxas-options=-v"])
func = mod.get_function("exponential_synapse")
func.prepare( [ np.int32, # syn_num
np.float64, # dt
np.intp, # spike list
np.intp, # pre-synaptic neuron list
np.intp, # tau; time constant
np.intp, # a; bump size
np.intp, # gmax array
np.intp, # eff; efficacy
np.intp ] ) # cond array
func.prepare('idPPPPPPP')
# [ np.int32, # syn_num
# np.float64, # dt
# np.intp, # spike list
# np.intp, # pre-synaptic neuron list
# np.intp, # tau; time constant
# np.intp, # a; bump size
# np.intp, # gmax array
# np.intp, # eff; efficacy
# np.intp ] ) # cond array
return func

def _get_update_I_non_cond_func(self):
mod = SourceModule(\
cuda_src_synapse_update_I % {"num": self.num},
options = ["--ptxas-options=-v"])
func = mod.get_function("get_input")
func.prepare([np.intp, # synapse state
np.intp, # cumulative dendrites number
np.intp, # dendrites number
np.intp, # pre-synaptic number ID
np.intp]) # output
func.prepare('PPPPP')
# [np.intp, # synapse state
# np.intp, # cumulative dendrites number
# np.intp, # dendrites number
# np.intp, # pre-synaptic number ID
# np.intp]) # output

self._block_get_input = (32,32,1)
self._grid_get_input = ((self.num - 1) / 32 + 1, 1)
Expand Down
4 changes: 3 additions & 1 deletion neurokernel/LPU/synapses/power_gpot_gpot.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -69,7 +69,9 @@ def get_update_func(self):
#Used 14 registers, 64 bytes cmem[0], 4 bytes cmem[16]
mod = SourceModule(template % {"n_synapse": self.num_synapse}, options = ["--ptxas-options=-v"])
func = mod.get_function("update_gpot_terminal_synapse")
func.prepare([np.intp, np.int32, np.int32, np.int32, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PiiiPPPPPPP')
#[np.intp, np.int32, np.int32, np.int32, np.intp,
# np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
self.block = (256,1,1)
self.grid = (min(6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, (self.num_synapse-1) / 256 + 1), 1)
return func
4 changes: 3 additions & 1 deletion neurokernel/LPU/synapses/power_gpot_gpot_sig.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,9 @@ def get_update_func(self):
#Used 14 registers, 64 bytes cmem[0], 4 bytes cmem[16]
mod = SourceModule(template % {"n_synapse": self.num_synapse}, options = ["--ptxas-options=-v"])
func = mod.get_function("update_gpot_terminal_synapse")
func.prepare([np.intp, np.int32, np.int32, np.int32, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
func.prepare('PiiiPPPPPPP')
#[np.intp, np.int32, np.int32, np.int32, np.intp, np.intp,
# np.intp, np.intp, np.intp, np.intp, np.intp])
self.block = (256,1,1)
self.grid = (min(6 * cuda.Context.get_device().MULTIPROCESSOR_COUNT, (self.num_synapse-1) / 256 + 1), 1)
return func
2 changes: 1 addition & 1 deletion neurokernel/LPU/utils/curand.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,7 @@ def get_curand_int_func():
"""
mod = SourceModule(code, no_extern_c = True)
func = mod.get_function("rand_setup")
func.prepare([np.intp, np.int32, np.uint64])
func.prepare('PiL')#[np.intp, np.int32, np.uint64])
return func


0 comments on commit a4c80c8

Please sign in to comment.