def _generate_runtime(self):
h_gen = CGenerator('{}/code_gen/runtime.h'.format(pkg_dir))
h_gen.if_define('runtime.h')
h_gen.blank_line(2)
h_gen.include("connection.h")
h_gen.blank_line(2)
h_gen.block('const int MAX_DELAY = {};'.format(self.max_delay))
h_gen.block('const int MIN_DELAY = {};'.format(self.min_delay))
h_gen.block('const {} G_MAX = {};'.format(real, self.g_max))
h_gen.block('const {} G_MIN = {};'.format(real, self.g_min))
for i in range(len(self.neuron_models)):
block_size = 32
h_gen.block("const int {}_BLOCK_SIZE = {};".format(
self.neuron_models[i].name.upper(), int(block_size)))
h_gen.block("const int {}_GRID_SIZE = {};".format(
self.neuron_models[i].name.upper(),
math.ceil(self.neuron_nums[i + 1] / block_size)))
for i in range(len(self.synapse_models)):
block_size = 128
h_gen.block("const int {}_BLOCK_SIZE = {};".format(
self.synapse_models[i].name.upper(), int(block_size)))
h_gen.block("const int {}_GRID_SIZE = {};".format(
self.synapse_models[i].name.upper(),
math.ceil(self.synapse_nums[i + 1] / block_size)))
h_gen.blank_line()
h_gen.block("const int MAX_BLOCK_SIZE = {};".format(
self.max_block_size))
h_gen.block("const int FIRED_TABLE_SIZE = {};".format(self.neuron_num))
h_gen.block('extern __device__ int * g_fired_table;')
h_gen.block('extern __device__ int * g_fired_table_sizes;')
for model in self.neuron_models:
h_gen.block('extern __device__ int * g_active_{}_table;'.format(
model.name.lower()))
h_gen.block('extern __device__ int g_active_{}_table_size;'.format(
model.name.lower()))
for model in self.synapse_models:
h_gen.block(
'extern __device__ CConnection * g_connection_{};'.format(
model.name.lower()))
external = set()
for model in self.neuron_models:
external |= set(model.parameters['external'])
for model in self.synapse_models:
external |= set(model.parameters['external'])
external -= set('t')
for i in external:
h_gen.block('extern __device__ {} * {};'.format(real, i))
h_gen.blank_line()
h_gen.block('extern "C" {')
h_gen.block('\tvoid **init_runtime(CConnection **connections);')
h_gen.block('}')
h_gen.blank_line()
h_gen.block(
'__device__ int commit2globalTable(int *shared_buf, volatile unsigned int size, '
'int *global_buf, int * global_size, int offset);')
h_gen.end_if_define('runtime.h')
h_gen.close()
cu_gen = CUDAGenerator('{}/code_gen/runtime.cu'.format(pkg_dir))
cu_gen.include_std("stdio.h")
cu_gen.include('../c_code/helper_cuda.h')
cu_gen.include("runtime.h")
cu_gen.blank_line(2)
cu_gen.block('__device__ int * g_fired_table;')
cu_gen.block('__device__ int * g_fired_table_sizes;')
for model in self.neuron_models:
cu_gen.block('__device__ int * g_active_{}_table;'.format(
model.name.lower()))
cu_gen.block('__device__ int g_active_{}_table_size;'.format(
model.name.lower()))
for model in self.synapse_models:
cu_gen.block('__device__ CConnection * g_connection_{};'.format(
model.name.lower()))
external = set()
for model in self.neuron_models:
external |= set(model.parameters['external'])
for model in self.synapse_models:
external |= set(model.parameters['external'])
external -= set('t')
for i in external:
cu_gen.block('__device__ {} * {};'.format(real, i))
cu_gen.blank_line(2)
cu_gen.block('void **init_runtime(CConnection ** connections)')
cu_gen.block('{')
cu_gen.block('\tint zero = 0;')
cu_gen.block('\tint *p_int = NULL;')
cu_gen.block('\t{} *p_{} = NULL;'.format(real, real))
cu_gen.blank_line()
cu_gen.block(
'\tvoid **ret = static_cast<void**>(malloc(sizeof(void*) * {}));'.
format(2))
cu_gen.blank_line()
cu_gen.malloc_symbol(symbol='g_fired_table_sizes',
gpu='p_int',
type_='int',
num='{}'.format(self.max_delay + 1))
cu_gen.block('\tret[0] = static_cast<void*>(p_int);')
cu_gen.malloc_symbol(symbol='g_fired_table',
gpu='p_int',
type_='int',
num='{}'.format(self.neuron_num *
(self.max_delay + 1)))
cu_gen.block('\tret[1] = static_cast<void*>(p_int);')
# cu_gen.block('\tprintf("\\n%p, %p, %p\\n", ret, ret[0], ret[1]);')
cu_gen.blank_line()
for model in self.neuron_models:
cu_gen.cu_line(
'cudaMemcpyToSymbol(g_active_{}_table_size, &zero, sizeof(int))'
.format(model.name.lower()))
cu_gen.malloc_symbol(symbol='g_active_{}_table'.format(
model.name.lower()),
gpu='p_int',
type_='int',
num='{}'.format(self.neuron_num))
cu_gen.block('\n')
for i in external:
cu_gen.malloc_symbol(symbol='{}'.format(i),
gpu='p_{}'.format(real),
type_=real,
num='{}'.format(self.neuron_num))
for i, model in enumerate(self.synapse_models):
cu_gen.cu_line(
'cudaMemcpyToSymbol(g_connection_{}, &(connections[{}]), sizeof(CConnection*))'
.format(model.name.lower(), i))
cu_gen.block('\treturn ret;')
cu_gen.block('}')
cu_gen.close()
def _generate_data_cu(self):
cu_gen = CUDAGenerator(
'{}/code_gen/connection.data.cu'.format(pkg_dir))
cu_gen.blank_line(2)
if self.debug:
cu_gen.include_std('stdio.h')
cu_gen.include_std('stdlib.h')
cu_gen.blank_line()
cu_gen.include('../c_code/helper_cuda.h')
cu_gen.include('connection.h')
cu_gen.blank_line(2)
cu_gen.line_no_end("CConnection * to_gpu_connection(CConnection *cpu)",
0)
cu_gen.open_brace()
cu_gen.line(
'CConnection * gpu = (CConnection*)malloc(sizeof(CConnection))')
cu_gen.line('gpu->n_len = cpu->n_len')
cu_gen.line('gpu->r_n_len = cpu->r_n_len')
cu_gen.line('gpu->s_len = cpu->s_len')
cu_gen.to_gpu(ret='gpu->delay_start',
cpu='cpu->delay_start',
type_='int',
num='cpu->n_len')
cu_gen.to_gpu(ret='gpu->delay_num',
cpu='cpu->delay_num',
type_='int',
num='cpu->n_len')
cu_gen.to_gpu(ret='gpu->rev_delay_start',
cpu='cpu->rev_delay_start',
type_='int',
num='cpu->r_n_len')
cu_gen.to_gpu(ret='gpu->rev_delay_num',
cpu='cpu->rev_delay_num',
type_='int',
num='cpu->r_n_len')
cu_gen.to_gpu(ret='gpu->rev_map2sid',
cpu='cpu->rev_map2sid',
type_='int',
num='cpu->s_len')
cu_gen.line('CConnection * ret = NULL')
cu_gen.to_gpu(ret='ret', cpu='gpu', type_='CConnection')
if self.debug:
cu_gen.line(line=r'printf("GPU CConnection Pointer: %p\n", ret)')
cu_gen.line(
line=
r'printf("GPU n_len: %d r_n_len: %d s_len: %d\n", gpu->n_len, gpu->r_n_len, gpu->s_len)'
)
cu_gen.line('return ret')
cu_gen.close_brace()
cu_gen.blank_line()
cu_gen.line_no_end(
"CConnection * from_gpu_connection(CConnection *gpu)", 0)
cu_gen.open_brace()
cu_gen.from_gpu(gpu='gpu', ret='ret', type_='CConnection')
if self.debug:
cu_gen.line(line=r'printf("CPU CConnection Pointer: %p\n", ret)')
cu_gen.line(
line=
r'printf("CPU n_len: %d r_n_len: %d s_len: %d\n", ret->n_len, ret->r_n_len, ret->s_len)'
)
cu_gen.line('return ret')
cu_gen.close_brace()
cu_gen.blank_line()
cu_gen.close()
return
def generate_data_cu(self, debug=False):
cu_gen = CUDAGenerator('{}/code_gen/{}.data.cu'.format(
pkg_dir, self.name.lower()))
cu_gen.blank_line(2)
cu_gen.include_std('stdlib.h')
cu_gen.blank_line()
cu_gen.include("../c_code/helper_cuda.h")
cu_gen.include("{}.h".format(self.name.lower()))
cu_gen.blank_line(2)
cu_gen.line_no_end(
"{} * to_gpu_{}({} *cpu, int num)".format(self.name.capitalize(),
self.name.lower(),
self.name.capitalize()),
0)
cu_gen.open_brace()
cu_gen.malloc(ret='gpu', type_=self.name.capitalize())
for i in self.parameters['special']:
if i != 'delay':
cu_gen.to_gpu(ret='gpu->p_{}'.format(i),
cpu='cpu->p_{}'.format(i),
num='num',
type_='int')
for i in self.parameters['constant']:
cu_gen.to_gpu(ret='gpu->p_{}'.format(i),
cpu='cpu->p_{}'.format(i),
num='num',
type_=real)
for i in self.parameters['variable']:
cu_gen.to_gpu(ret='gpu->p_{}'.format(i),
cpu='cpu->p_{}'.format(i),
num='num',
type_=real)
cu_gen.line('{} * ret = NULL'.format(self.name.capitalize()))
cu_gen.to_gpu(ret='ret',
cpu='gpu',
num='1',
type_=self.name.capitalize())
cu_gen.line('return ret')
cu_gen.close_brace()
cu_gen.blank_line()
cu_gen.line_no_end(
"{} * from_gpu_{}({} *gpu, int num)".format(
self.name.capitalize(), self.name.lower(),
self.name.capitalize()), 0)
cu_gen.open_brace()
cu_gen.from_gpu(ret='cpu', gpu='gpu', type_=self.name.capitalize())
cu_gen.line('return cpu')
cu_gen.close_brace()
cu_gen.blank_line()
cu_gen.close()
return