Skip to content

Commit

Permalink
baseneuron to check if self.I exists, hard coded constants in kernel
Browse files Browse the repository at this point in the history
  • Loading branch information
@yiyin
yiyin committed Dec 10, 2015
1 parent a1a6811 commit 7d1bc90
Showing 1 changed file with 36 additions and 21 deletions.
57 changes: 36 additions & 21 deletions neurokernel/LPU/neurons/baseneuron.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
Base neuron class used by LPU.
"""

import warnings
from abc import ABCMeta, abstractmethod, abstractproperty
import os.path
import numpy as np
Expand Down Expand Up @@ -68,6 +69,8 @@ def __init__(self, n_dict, neuron_state_pointer, dt, debug, LPU_id=None):
debug is a boolean and is intended to be used for debugging purposes.
'''
self.__LPU_id = LPU_id

self.__neuron_state_pointer = neuron_state_pointer
self.__num_neurons = len(n_dict['id'])
_num_dendrite_cond = np.asarray([n_dict['num_dendrites_cond'][i]
Expand All @@ -90,10 +93,20 @@ def __init__(self, n_dict, neuron_state_pointer, dt, debug, LPU_id=None):
dtype=np.int32))
self.__V_rev = garray.to_gpu(np.asarray(n_dict['reverse'],
dtype=np.double))
self.I = garray.zeros(self.__num_neurons, np.double)
try:
assert(isinstance(getattr(self, 'I'), garray.GPUArray))
except AttributeError, AssertionError:
self.I = garray.zeros(self.__num_neurons, np.double)

if self.I.size != self.__num_neurons:
self.I = garray.zeros(self.__num_neurons, np.double)
warnings.warn('GPUArray representing current in ' +
'{}'.format(type(self).__name__) + ' neuron of ' +
'{}'.format(self.__LPU_id) + ' LPU has been changed ' +
' due to incompatible size.')

self.__update_I_cond = self.__get_update_I_cond_func()
self.__update_I_non_cond = self.__get_update_I_non_cond_func()
self.__LPU_id = LPU_id
self.__debug = debug
if self.__debug:
if self.__LPU_id is None:
Expand Down Expand Up @@ -149,13 +162,13 @@ def update_I(self, synapse_state, st=None, logger=None):
self.I.fill(0)
if self.__pre.size > 0:
self.__update_I_non_cond.prepared_async_call(
self.__grid_get_input, self.__block_get_input, st,
self.__grid_get_input_I, self.__block_get_input_I, st,
int(synapse_state), self.__cum_num_dendrite.gpudata,
self.__num_dendrite.gpudata, self.__pre.gpudata,
self.I.gpudata)
if self.__cond_pre.size > 0:
self.__update_I_cond.prepared_async_call(
self.__grid_get_input, self.__block_get_input, st,
self.__grid_get_input_cond, self.__block_get_input_cond, st,
int(synapse_state), self.__cum_num_dendrite_cond.gpudata,
self.__num_dendrite_cond.gpudata, self.__cond_pre.gpudata,
self.I.gpudata, int(self.__neuron_state_pointer),
Expand All @@ -180,13 +193,13 @@ def __post_run(self):

def __get_update_I_cond_func(self):
template = """
#define N 32
#define NUM_NEURONS %(num_neurons)d
__global__ void get_input(double* synapse, int* cum_num_dendrite,
int* num_dendrite, int* pre, double* I_pre,
double* V, double* V_rev)
{
// must use block size (32, 32, 1)
int tidx = threadIdx.x;
int tidy = threadIdx.y;
int bid = blockIdx.x;
Expand All @@ -200,15 +213,15 @@ def __get_update_I_cond_func(self):
if(tidy == 0)
{
neuron = bid * N + tidx;
neuron = bid * 32 + tidx;
if(neuron < NUM_NEURONS)
{
num_den[tidx] = num_dendrite[neuron];
V_in[tidx] = V[neuron];
}
} else if(tidy == 1)
{
neuron = bid * N + tidx;
neuron = bid * 32 + tidx;
if(neuron < NUM_NEURONS)
{
den_start[tidx] = cum_num_dendrite[neuron];
Expand All @@ -219,15 +232,15 @@ def __get_update_I_cond_func(self):
__syncthreads();
neuron = bid * N + tidy ;
neuron = bid * 32 + tidy ;
if(neuron < NUM_NEURONS)
{
int n_den = num_den[tidy];
int start = den_start[tidy];
double VV = V_in[tidy];
for(int i = tidx; i < n_den; i += N)
for(int i = tidx; i < n_den; i += 32)
{
input[tidy][tidx] += synapse[pre[start + i]] * (VV - V_rev[start + i]);
}
Expand Down Expand Up @@ -261,7 +274,7 @@ def __get_update_I_cond_func(self):
if(tidy == 0)
{
input[tidx][0] += input[tidx][1];
neuron = bid*N+tidx;
neuron = bid*32+tidx;
if(neuron < NUM_NEURONS)
{
I_pre[neuron] -= input[tidx][0];
Expand All @@ -273,20 +286,20 @@ 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])
self.__block_get_input = (32, 32, 1)
self.__grid_get_input = ((self.__num_neurons - 1) / 32 + 1, 1)
func.prepare('PPPPPPP')
#[np.intp, np.intp, np.intp, np.intp, np.intp, np.intp, np.intp])
self.__block_get_input_cond = (32, 32, 1)
self.__grid_get_input_cond = ((self.__num_neurons - 1) / 32 + 1, 1)
return func

def __get_update_I_non_cond_func(self):
template = """
#define N 32
#define NUM_NEURONS %(num_neurons)d
__global__ void get_input(double* synapse, int* cum_num_dendrite,
int* num_dendrite, int* pre, double* I_pre)
{
// must use block size (32, 32, 1)
int tidx = threadIdx.x;
int tidy = threadIdx.y;
int bid = blockIdx.x;
Expand All @@ -299,14 +312,14 @@ def __get_update_I_non_cond_func(self):
if(tidy == 0)
{
neuron = bid * N + tidx;
neuron = bid * 32 + tidx;
if(neuron < NUM_NEURONS)
{
num_den[tidx] = num_dendrite[neuron];
}
} else if(tidy == 1)
{
neuron = bid * N + tidx;
neuron = bid * 32 + tidx;
if(neuron < NUM_NEURONS)
{
den_start[tidx] = cum_num_dendrite[neuron];
Expand All @@ -317,13 +330,13 @@ def __get_update_I_non_cond_func(self):
__syncthreads();
neuron = bid * N + tidy ;
neuron = bid * 32 + tidy ;
if(neuron < NUM_NEURONS){
int n_den = num_den[tidy];
int start = den_start[tidy];
for(int i = tidx; i < n_den; i += N)
for(int i = tidx; i < n_den; i += 32)
{
input[tidy][tidx] += synapse[pre[start] + i];
}
Expand Down Expand Up @@ -356,7 +369,7 @@ def __get_update_I_non_cond_func(self):
if(tidy == 0)
{
input[tidx][0] += input[tidx][1];
neuron = bid*N+tidx;
neuron = bid*32+tidx;
if(neuron < NUM_NEURONS)
{
I_pre[neuron] += input[tidx][0];
Expand All @@ -369,5 +382,7 @@ 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])
self.__block_get_input_I = (32, 32, 1)
self.__grid_get_input_I = ((self.__num_neurons - 1) / 32 + 1, 1)
return func

0 comments on commit 7d1bc90

Please sign in to comment.