From 35496b6a648d98dc41d6226c1d43650aba58cdfc Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Mon, 14 Mar 2016 16:10:59 +0800 Subject: add seq_buffer --- nerv/examples/network_debug/config.lua | 2 +- nerv/examples/network_debug/main.lua | 19 ++---- nerv/examples/network_debug/network.lua | 114 +++++++++++++++----------------- nerv/examples/network_debug/reader.lua | 78 +++++----------------- nerv/io/seq_buffer.lua | 105 +++++++++++++++++++++++++++++ 5 files changed, 182 insertions(+), 136 deletions(-) diff --git a/nerv/examples/network_debug/config.lua b/nerv/examples/network_debug/config.lua index e20d5a9..9025b78 100644 --- a/nerv/examples/network_debug/config.lua +++ b/nerv/examples/network_debug/config.lua @@ -12,7 +12,7 @@ function get_global_conf() layer_num = 1, chunk_size = 15, batch_size = 20, - max_iter = 35, + max_iter = 3, param_random = function() return (math.random() / 5 - 0.1) end, dropout_rate = 0.5, timer = nerv.Timer(), diff --git a/nerv/examples/network_debug/main.lua b/nerv/examples/network_debug/main.lua index 790c404..1bee43c 100644 --- a/nerv/examples/network_debug/main.lua +++ b/nerv/examples/network_debug/main.lua @@ -6,29 +6,20 @@ nerv.include(arg[1]) local global_conf = get_global_conf() local timer = global_conf.timer -timer:tic('IO') - local data_path = 'examples/lmptb/PTBdata/' -local train_reader = nerv.Reader(data_path .. 'vocab', data_path .. 'ptb.train.txt.adds') -local val_reader = nerv.Reader(data_path .. 'vocab', data_path .. 'ptb.valid.txt.adds') - -local train_data = train_reader:get_all_batch(global_conf) -local val_data = val_reader:get_all_batch(global_conf) local layers = get_layers(global_conf) local connections = get_connections(global_conf) -local NN = nerv.NN(global_conf, train_data, val_data, layers, connections) - -timer:toc('IO') -timer:check('IO') -io.flush() +local NN = nerv.NN(global_conf, layers, connections) timer:tic('global') local best_cv = 1e10 for i = 1, global_conf.max_iter do timer:tic('Epoch' .. i) - local train_ppl, val_ppl = NN:epoch() + local train_reader = nerv.Reader(data_path .. 'vocab', data_path .. 'ptb.train.txt.adds') + local val_reader = nerv.Reader(data_path .. 'vocab', data_path .. 'ptb.valid.txt.adds') + local train_ppl, val_ppl = NN:epoch(train_reader, val_reader) if val_ppl < best_cv then best_cv = val_ppl else @@ -43,3 +34,5 @@ timer:toc('global') timer:check('global') timer:check('network') timer:check('gc') +timer:check('IO') +global_conf.cumat_type.print_profile() diff --git a/nerv/examples/network_debug/network.lua b/nerv/examples/network_debug/network.lua index 5518e27..1841d21 100644 --- a/nerv/examples/network_debug/network.lua +++ b/nerv/examples/network_debug/network.lua @@ -2,11 +2,17 @@ nerv.include('select_linear.lua') local nn = nerv.class('nerv.NN') -function nn:__init(global_conf, train_data, val_data, layers, connections) +function nn:__init(global_conf, layers, connections) self.gconf = global_conf self.network = self:get_network(layers, connections) - self.train_data = self:get_data(train_data) - self.val_data = self:get_data(val_data) + + self.output = {} + self.err_output = {} + for i = 1, self.gconf.chunk_size do + self.output[i] = {self.gconf.cumat_type(self.gconf.batch_size, 1)} + self.err_output[i] = {self.gconf.cumat_type(self.gconf.batch_size, 1)} + self.err_output[i][2] = self.gconf.cumat_type(self.gconf.batch_size, 1) + end end function nn:get_network(layers, connections) @@ -20,79 +26,67 @@ function nn:get_network(layers, connections) return network end -function nn:get_data(data) - local err_output = {} - local softmax_output = {} - local output = {} - for i = 1, self.gconf.chunk_size do - err_output[i] = self.gconf.cumat_type(self.gconf.batch_size, 1) - softmax_output[i] = self.gconf.cumat_type(self.gconf.batch_size, self.gconf.vocab_size) - output[i] = self.gconf.cumat_type(self.gconf.batch_size, 1) - end - local ret = {} - for i = 1, #data do - ret[i] = {} - ret[i].input = {} - ret[i].output = {} - ret[i].err_input = {} - ret[i].err_output = {} - for t = 1, self.gconf.chunk_size do - ret[i].input[t] = {} - ret[i].output[t] = {} - ret[i].err_input[t] = {} - ret[i].err_output[t] = {} - ret[i].input[t][1] = data[i].input[t] - ret[i].input[t][2] = data[i].output[t] - ret[i].output[t][1] = output[t] - local err_input = self.gconf.mmat_type(self.gconf.batch_size, 1) - for j = 1, self.gconf.batch_size do - if t <= data[i].seq_len[j] then - err_input[j - 1][0] = 1 - else - err_input[j - 1][0] = 0 +function nn:process(data, do_train, reader) + local timer = self.gconf.timer + local buffer = nerv.SeqBuffer(self.gconf, { + batch_size = self.gconf.batch_size, chunk_size = self.gconf.chunk_size, + readers = {reader}, + }) + local total_err = 0 + local total_frame = 0 + self.network:epoch_init() + while true do + timer:tic('IO') + data = buffer:get_data() + if data == nil then + break + end + local err_input = {} + if do_train then + for t = 1, self.gconf.chunk_size do + local tmp = self.gconf.mmat_type(self.gconf.batch_size, 1) + for i = 1, self.gconf.batch_size do + if t <= data.seq_length[i] then + tmp[i - 1][0] = 1 + else + tmp[i - 1][0] = 0 + end end + err_input[t] = {self.gconf.cumat_type.new_from_host(tmp)} end - ret[i].err_input[t][1] = self.gconf.cumat_type.new_from_host(err_input) - ret[i].err_output[t][1] = err_output[t] - ret[i].err_output[t][2] = softmax_output[t] end - ret[i].seq_length = data[i].seq_len - ret[i].new_seq = {} - for j = 1, self.gconf.batch_size do - if data[i].seq_start[j] then - table.insert(ret[i].new_seq, j) - end + local info = {input = {}, output = self.output, err_input = err_input, do_train = do_train, + err_output = self.err_output, seq_length = data.seq_length, new_seq = data.new_seq} + for t = 1, self.gconf.chunk_size do + info.input[t] = {data.data['input'][t]} + info.input[t][2] = data.data['label'][t] end - end - return ret -end + timer:toc('IO') -function nn:process(data, do_train) - local timer = self.gconf.timer - local total_err = 0 - local total_frame = 0 - self.network:epoch_init() - for id = 1, #data do - data[id].do_train = do_train timer:tic('network') - self.network:mini_batch_init(data[id]) + self.network:mini_batch_init(info) self.network:propagate() timer:toc('network') + + timer:tic('IO') for t = 1, self.gconf.chunk_size do - local tmp = data[id].output[t][1]:new_to_host() + local tmp = info.output[t][1]:new_to_host() for i = 1, self.gconf.batch_size do - if t <= data[id].seq_length[i] then + if t <= info.seq_length[i] then total_err = total_err + math.log10(math.exp(tmp[i - 1][0])) total_frame = total_frame + 1 end end end + timer:toc('IO') + + timer:tic('network') if do_train then - timer:tic('network') self.network:back_propagate() self.network:update() - timer:toc('network') end + timer:toc('network') + timer:tic('gc') collectgarbage('collect') timer:toc('gc') @@ -100,11 +94,11 @@ function nn:process(data, do_train) return math.pow(10, - total_err / total_frame) end -function nn:epoch() - local train_error = self:process(self.train_data, true) +function nn:epoch(train_reader, val_reader) + local train_error = self:process(self.train_data, true, train_reader) local tmp = self.gconf.dropout_rate self.gconf.dropout_rate = 0 - local val_error = self:process(self.val_data, false) + local val_error = self:process(self.val_data, false, val_reader) self.gconf.dropout_rate = tmp return train_error, val_error end diff --git a/nerv/examples/network_debug/reader.lua b/nerv/examples/network_debug/reader.lua index b10baaf..70c0c97 100644 --- a/nerv/examples/network_debug/reader.lua +++ b/nerv/examples/network_debug/reader.lua @@ -3,6 +3,7 @@ local Reader = nerv.class('nerv.Reader') function Reader:__init(vocab_file, input_file) self:get_vocab(vocab_file) self:get_seq(input_file) + self.offset = 1 end function Reader:get_vocab(vocab_file) @@ -31,7 +32,8 @@ end function Reader:get_seq(input_file) local f = io.open(input_file, 'r') self.seq = {} - while true do + -- while true do + for i = 1, 26 do local seq = f:read() if seq == nil then break @@ -47,67 +49,19 @@ function Reader:get_seq(input_file) end end -function Reader:get_in_out(id, pos) - return self.seq[id][pos], self.seq[id][pos + 1], pos + 1 == #self.seq[id] -end - -function Reader:get_all_batch(global_conf) - local data = {} - local pos = {} - local offset = 1 - for i = 1, global_conf.batch_size do - pos[i] = nil +function Reader:get_data() + if self.offset > #self.seq then + return nil end - while true do - -- for i = 1, 26 do - local input = {} - local output = {} - for i = 1, global_conf.chunk_size do - input[i] = global_conf.mmat_type(global_conf.batch_size, 1) - input[i]:fill(global_conf.nn_act_default) - output[i] = global_conf.mmat_type(global_conf.batch_size, 1) - output[i]:fill(global_conf.nn_act_default) - end - local seq_start = {} - local seq_end = {} - local seq_len = {} - for i = 1, global_conf.batch_size do - seq_start[i] = false - seq_end[i] = false - seq_len[i] = 0 - end - local has_new = false - for i = 1, global_conf.batch_size do - if pos[i] == nil then - if offset < #self.seq then - seq_start[i] = true - pos[i] = {offset, 1} - offset = offset + 1 - end - end - if pos[i] ~= nil then - has_new = true - for j = 1, global_conf.chunk_size do - local final - input[j][i-1][0], output[j][i-1][0], final = self:get_in_out(pos[i][1], pos[i][2]) - seq_len[i] = j - if final then - seq_end[i] = true - pos[i] = nil - break - end - pos[i][2] = pos[i][2] + 1 - end - end - end - if not has_new then - break - end - for i = 1, global_conf.chunk_size do - input[i] = global_conf.cumat_type.new_from_host(input[i]) - output[i] = global_conf.cumat_type.new_from_host(output[i]) - end - table.insert(data, {input = input, output = output, seq_start = seq_start, seq_end = seq_end, seq_len = seq_len}) + local tmp = self.seq[self.offset] + local res = { + input = nerv.MMatrixFloat(#tmp - 1, 1), + label = nerv.MMatrixFloat(#tmp - 1, 1), + } + for i = 1, #tmp - 1 do + res.input[i - 1][0] = tmp[i] + res.label[i - 1][0] = tmp[i + 1] end - return data + self.offset = self.offset + 1 + return res end diff --git a/nerv/io/seq_buffer.lua b/nerv/io/seq_buffer.lua index e69de29..ad1b3f7 100644 --- a/nerv/io/seq_buffer.lua +++ b/nerv/io/seq_buffer.lua @@ -0,0 +1,105 @@ +local SeqBuffer = nerv.class('nerv.SeqBuffer', 'nerv.DataBuffer') + +function SeqBuffer:__init(global_conf, buffer_conf) + self.gconf = global_conf + + self.batch_size = buffer_conf.batch_size + self.chunk_size = buffer_conf.chunk_size + self.readers = buffer_conf.readers + self.nn_act_default = buffer_conf.nn_act_default + if self.nn_act_default == nil then + self.nn_act_default = 0 + end + + self.mat_type = self.gconf.mmat_type + self.queue = {} + self.head = 1 + self.tail = 0 +end + +function SeqBuffer:new_mini_batch() + local res = {} + res.data = {} + res.new_seq = {} + res.seq_length = {} + for i = 1, self.batch_size do + res.seq_length[i] = 0 + end + return res +end + +function SeqBuffer:saturate(batch) + if self.queue[self.head] ~= nil and self.queue[self.head].seq_length[batch] ~= 0 then + return true + end + local data = {} + local drow = nil + for i = 1, #self.readers do + local tmp = self.readers[i]:get_data() + if tmp == nil then + return false + end + for id, d in pairs(tmp) do + if drow == nil then + drow = d:nrow() + elseif d:nrow() ~= drow then + nerv.error('readers provides with inconsistent rows of data') + end + data[id] = d + end + end + local offset = 0 + local head = self.head + while offset < drow do + local last = math.min(offset + self.chunk_size, drow) + if head > self.tail then + self.tail = self.tail + 1 + self.queue[self.tail] = self:new_mini_batch() + end + self.queue[head].seq_length[batch] = last - offset + if offset == 0 then + table.insert(self.queue[head].new_seq, batch) + end + local mini_batch = self.queue[head].data + for id, d in pairs(data) do + if mini_batch[id] == nil then + mini_batch[id] = {} + end + local tmp = mini_batch[id] + for i = offset + 1, last do + local chunk = i - offset + if tmp[chunk] == nil then + tmp[chunk] = self.mat_type(self.batch_size, d:ncol()) + tmp[chunk]:fill(self.nn_act_default) + end + tmp[chunk]:copy_from(d, i - 1, i, batch - 1) + end + end + head = head + 1 + offset = last + end + return true +end + +function SeqBuffer:get_data() + local has_data = false + for i = 1, self.batch_size do + if self:saturate(i) then + has_data = true + end + end + if not has_data then + return nil + end + local res = self.queue[self.head] + self.queue[self.head] = nil + self.head = self.head + 1 + if not self.gconf.use_cpu then + for id, d in pairs(res.data) do + for i = 1, #d do + d[i] = self.gconf.cumat_type.new_from_host(d[i]) + end + end + end + return res +end -- cgit v1.2.3 From b08da1fef90e93b188704056cdae651d7865f98d Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Mon, 14 Mar 2016 20:07:15 +0800 Subject: speedup border flush --- Makefile | 5 +- nerv/examples/network_debug/config.lua | 2 +- nerv/examples/network_debug/main.lua | 2 +- nerv/examples/network_debug/reader.lua | 4 +- nerv/nn/network.lua | 215 ++++++++++++++++++++------------- 5 files changed, 141 insertions(+), 87 deletions(-) diff --git a/Makefile b/Makefile index 28012da..3995d21 100644 --- a/Makefile +++ b/Makefile @@ -1,4 +1,4 @@ -.PHONY: all clean install luajit luarocks speech +.PHONY: all clean install luajit luarocks speech submodule ############## EDIT THESE LINES ##################### SHELL := /bin/bash PREFIX := $(CURDIR)/install/ @@ -26,7 +26,8 @@ export BLAS_LDFLAGS nerv-clean speech/speech_utils-clean speech/htk_io-clean speech/kaldi_io-clean speech/kaldi_decode-clean \ Penlight -all: luajit luarocks Penlight nerv +all: nerv +submodule: luajit luajit Penlight luajit: PREFIX=$(PREFIX) ./tools/build_luajit.sh luarocks: diff --git a/nerv/examples/network_debug/config.lua b/nerv/examples/network_debug/config.lua index 9025b78..093bde2 100644 --- a/nerv/examples/network_debug/config.lua +++ b/nerv/examples/network_debug/config.lua @@ -12,7 +12,7 @@ function get_global_conf() layer_num = 1, chunk_size = 15, batch_size = 20, - max_iter = 3, + max_iter = 1, param_random = function() return (math.random() / 5 - 0.1) end, dropout_rate = 0.5, timer = nerv.Timer(), diff --git a/nerv/examples/network_debug/main.lua b/nerv/examples/network_debug/main.lua index 1bee43c..bbcdb6c 100644 --- a/nerv/examples/network_debug/main.lua +++ b/nerv/examples/network_debug/main.lua @@ -20,12 +20,12 @@ for i = 1, global_conf.max_iter do local train_reader = nerv.Reader(data_path .. 'vocab', data_path .. 'ptb.train.txt.adds') local val_reader = nerv.Reader(data_path .. 'vocab', data_path .. 'ptb.valid.txt.adds') local train_ppl, val_ppl = NN:epoch(train_reader, val_reader) + nerv.printf('Epoch %d: %f %f %f\n', i, global_conf.lrate, train_ppl, val_ppl) if val_ppl < best_cv then best_cv = val_ppl else global_conf.lrate = global_conf.lrate / 2.0 end - nerv.printf('Epoch %d: %f %f %f\n', i, global_conf.lrate, train_ppl, val_ppl) timer:toc('Epoch' .. i) timer:check('Epoch' .. i) io.flush() diff --git a/nerv/examples/network_debug/reader.lua b/nerv/examples/network_debug/reader.lua index 70c0c97..76a78cf 100644 --- a/nerv/examples/network_debug/reader.lua +++ b/nerv/examples/network_debug/reader.lua @@ -32,8 +32,8 @@ end function Reader:get_seq(input_file) local f = io.open(input_file, 'r') self.seq = {} - -- while true do - for i = 1, 26 do + while true do + -- for i = 1, 26 do local seq = f:read() if seq == nil then break diff --git a/nerv/nn/network.lua b/nerv/nn/network.lua index 2cb83ce..910cdad 100644 --- a/nerv/nn/network.lua +++ b/nerv/nn/network.lua @@ -16,6 +16,7 @@ function network:__init(id, global_conf, network_conf) if self.nn_act_default == nil then self.nn_act_default = 0 end + self.layers = {} self.input_conn = {} self.output_conn = {} @@ -36,16 +37,41 @@ function network:__init(id, global_conf, network_conf) end self.output_conn[id][port] = {0, i, time} end + self.delay = 0 for i = 1, #self.layers do local dim_in, _ = self.layers[i]:get_dim() for j = 1, #dim_in do + if self.input_conn[i][j] == nil then + nerv.error('dangling input') + end local time = self.input_conn[i][j][3] if math.abs(time) > self.delay then self.delay = math.abs(time) end end end + + self.input_edge = {} + self.output_edge = {} + for t = -self.delay, self.delay do + self.input_edge[t] = {} + self.output_edge[t] = {} + end + for i = 1, #self.layers do + local dim_in, dim_out = self.layers[i]:get_dim() + for j = 1, #dim_in do + local time = self.input_conn[i][j][3] + table.insert(self.input_edge[time], {i, j}) + end + for j = 1, #dim_out do + if self.output_conn[i][j] == nil then + nerv.error('dangling output') + end + local time = self.output_conn[i][j][3] + table.insert(self.output_edge[time], {i, j}) + end + end end function network:compile(layer) @@ -112,11 +138,20 @@ function network:init(batch_size, chunk_size) self:make_initial_store() collectgarbage('collect') + + self.flush = {} + for t = 1, self.chunk_size do + self.flush[t] = {} + end end function network:epoch_init() + self.timestamp = 0 for i = 1, #self.layers do self.layers[i]:init(self.batch_size, self.chunk_size) + for t = 1, self.chunk_size do + self.flush[t][i] = {timestamp = 0, input = {}, output = {}} + end end end @@ -134,12 +169,10 @@ function network:topsort() for i = 1, #self.layers do local _, dim_out = self.layers[i]:get_dim() for j = 1, #dim_out do - if self.output_conn[i][j] ~= nil then - local edge = self.output_conn[i][j] - local id, time = edge[1], edge[3] + t - if time >= 1 and time <= self.chunk_size and id ~= 0 then - degree[time][id] = degree[time][id] + 1 - end + local edge = self.output_conn[i][j] + local id, time = edge[1], edge[3] + t + if time >= 1 and time <= self.chunk_size and id ~= 0 then + degree[time][id] = degree[time][id] + 1 end end end @@ -161,15 +194,13 @@ function network:topsort() l = l + 1 local _, dim_out = self.layers[i]:get_dim() for j = 1, #dim_out do - if self.output_conn[i][j] ~= nil then - local edge = self.output_conn[i][j] - local id, time = edge[1], edge[3] + t - if time >= 1 and time <= self.chunk_size and id ~= 0 then - degree[time][id] = degree[time][id] - 1 - if degree[time][id] == 0 then - r = r + 1 - self.queue[r] = {chunk = time, id = id} - end + local edge = self.output_conn[i][j] + local id, time = edge[1], edge[3] + t + if time >= 1 and time <= self.chunk_size and id ~= 0 then + degree[time][id] = degree[time][id] - 1 + if degree[time][id] == 0 then + r = r + 1 + self.queue[r] = {chunk = time, id = id} end end end @@ -202,17 +233,19 @@ function network:make_initial_store() memory[t][i][j]:fill(self.nn_act_default) end end - -- memory[t][0] stores network input - memory[t][0] = {} - for j = 1, #self.dim_in do - memory[t][0][j] = self.mat_type(self.batch_size, self.dim_in[j]) - memory[t][0][j]:fill(self.nn_act_default) - end - -- err_memory[t][0] stores network err_input - err_memory[t][0] = {} - for j = 1, #self.dim_out do - err_memory[t][0][j] = self.mat_type(self.batch_size, self.dim_out[j]) - err_memory[t][0][j]:fill(0) + if t < 1 or t > self.chunk_size then + -- memory[t][0] stores network input + memory[t][0] = {} + for j = 1, #self.dim_in do + memory[t][0][j] = self.mat_type(self.batch_size, self.dim_in[j]) + memory[t][0][j]:fill(self.nn_act_default) + end + -- err_memory[t][0] stores network err_input + err_memory[t][0] = {} + for j = 1, #self.dim_out do + err_memory[t][0][j] = self.mat_type(self.batch_size, self.dim_out[j]) + err_memory[t][0][j]:fill(0) + end end end @@ -314,9 +347,14 @@ function network:make_initial_store() self.legacy[t] = {} for i = 1, #self.layers do self.legacy[t][i] = {} - local _, dim_out = self.layers[i]:get_dim() - for j = 1, #dim_out do - self.legacy[t][i][j] = memory[t][i][j] + end + end + for d = 1, self.delay do + for t = 1 - d, 0 do + for i = 1, #self.output_edge[d] do + local edge = self.output_edge[d][i] + local id, port = edge[1], edge[2] + self.legacy[t][id][port] = memory[t][id][port] end end end @@ -383,59 +421,74 @@ function network:mini_batch_init(info) self.info = info self:set_input(self.info.input) self:set_output(self.info.output) + if self.info.do_train then + self:set_err_input(self.info.err_input) + self:set_err_output(self.info.err_output) + end -- calculate border self.max_length = 0 - self.border = {} - for i = 1, self.chunk_size do - self.border[i] = {} - end + self.timestamp = self.timestamp + 1 for i = 1, self.batch_size do if self.info.seq_length[i] > self.max_length then self.max_length = self.info.seq_length[i] end - for t = 1, self.delay do - local chunk = self.info.seq_length[i] + t - if chunk > self.chunk_size then - break + local border = self.info.seq_length[i] + for d = 1, self.delay do + for t = border + 1, border + d do + if t > self.chunk_size then + break + end + for j = 1, #self.output_edge[-d] do + local edge = self.output_edge[-d][j] + local id, port = edge[1], edge[2] + local flush = self.flush[t][id] + if flush.timestamp ~= self.timestamp then + flush.timestamp = self.timestamp + flush.input = {} + flush.output = {} + end + table.insert(flush.output, {port, i}) + end + end + if self.info.do_train then + for t = border, border - d + 1, -1 do + if t < 1 then + break + end + for j = 1, #self.input_edge[-d] do + local edge = self.input_edge[-d][j] + local id, port = edge[1], edge[2] + local flush = self.flush[t][id] + if flush.timestamp ~= self.timestamp then + flush.timestamp = self.timestamp + flush.input = {} + flush.output = {} + end + table.insert(flush.input, {port, i}) + end + end end - table.insert(self.border[chunk], i) end end -- copy legacy - for t = 1 - self.delay, 0 do - for i = 1, #self.layers do - local _, dim_out = self.layers[i]:get_dim() - for j = 1, #dim_out do - if t + self.chunk_size >= 1 and self.output_conn[i][j][1] ~= 0 then - self.legacy[t][i][j]:copy_from(self.output[t + self.chunk_size][i][j]) + for d = 1, self.delay do + for t = 1 - d, 0 do + for i = 1, #self.output_edge[d] do + local edge = self.output_edge[d][i] + local id, port = edge[1], edge[2] + if t + self.chunk_size >= 1 and self.output_conn[id][port][1] ~= 0 then + self.legacy[t][id][port]:copy_from(self.output[t + self.chunk_size][id][port]) end - for k = 1, #self.info.new_seq do - local batch = self.info.new_seq[k] - self.legacy[t][i][j][batch - 1]:fill(self.nn_act_default) + for j = 1, #self.info.new_seq do + local batch = self.info.new_seq[j] + self.legacy[t][id][port][batch - 1]:fill(self.nn_act_default) end end end end - if self.info.do_train then - self:set_err_input(self.info.err_input) - self:set_err_output(self.info.err_output) - - -- flush border gradient - for t = self.max_length + 1, self.max_length + self.delay do - if t > self.chunk_size then - break - end - for i = 1, #self.layers do - local dim_in, _ = self.layers[i]:get_dim() - for j = 1, #dim_in do - self.err_output[t][i][j]:fill(0) - end - end - end - end end function network:propagate() @@ -445,11 +498,11 @@ function network:propagate() self.layers[id]:propagate(self.input[t][id], self.output[t][id], t) end -- flush border activation - for j = 1, #self.border[t] do - local batch = self.border[t][j] - local _, dim_out = self.layers[id]:get_dim() - for k = 1, #dim_out do - self.output[t][id][k][batch - 1]:fill(self.nn_act_default) + if self.flush[t][id].timestamp == self.timestamp then + for j = 1, #self.flush[t][id].output do + local border = self.flush[t][id].output[j] + local port, batch = border[1], border[2] + self.output[t][id][port][batch - 1]:fill(self.nn_act_default) end end end @@ -459,15 +512,8 @@ function network:back_propagate() for i = #self.queue, 1, -1 do local t, id = self.queue[i].chunk, self.queue[i].id if t <= self.max_length then - -- flush border gradient - for j = 1, #self.border[t] do - local batch = self.border[t][j] - local _, dim_out = self.layers[id]:get_dim() - for k = 1, #dim_out do - self.err_input[t][id][k][batch - 1]:fill(0) - end - end self.layers[id]:back_propagate(self.err_input[t][id], self.err_output[t][id], self.input[t][id], self.output[t][id], t) + -- gradient clip if self.clip ~= nil then local dim_in, _ = self.layers[id]:get_dim() for j = 1, #dim_in do @@ -475,14 +521,21 @@ function network:back_propagate() end end end + -- flush border gradient + if self.flush[t][id].timestamp == self.timestamp then + for j = 1, #self.flush[t][id].input do + local border = self.flush[t][id].input[j] + local port, batch = border[1], border[2] + self.err_output[t][id][port][batch - 1]:fill(0) + end + end end end function network:update() - for i = 1, #self.queue do - local t, id = self.queue[i].chunk, self.queue[i].id - if t <= self.max_length then - self.layers[id]:update(self.err_input[t][id], self.input[t][id], self.output[t][id], t) + for t = 1, self.max_length do + for i = 1, #self.layers do + self.layers[i]:update(self.err_input[t][i], self.input[t][i], self.output[t][i], t) end end end -- cgit v1.2.3 From 51a3beef4a7cbd94278a406664212b6597aedd93 Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Tue, 15 Mar 2016 10:56:35 +0800 Subject: speedup duplicate layer --- Makefile | 2 +- nerv/examples/network_debug/config.lua | 2 +- nerv/layer/duplicate.lua | 5 +---- nerv/nn/network.lua | 30 +++++++++++++++++++++++++++--- 4 files changed, 30 insertions(+), 9 deletions(-) diff --git a/Makefile b/Makefile index 3995d21..9c9c86c 100644 --- a/Makefile +++ b/Makefile @@ -27,7 +27,7 @@ export BLAS_LDFLAGS Penlight all: nerv -submodule: luajit luajit Penlight +submodule: luajit luarocks Penlight luajit: PREFIX=$(PREFIX) ./tools/build_luajit.sh luarocks: diff --git a/nerv/examples/network_debug/config.lua b/nerv/examples/network_debug/config.lua index 093bde2..e20d5a9 100644 --- a/nerv/examples/network_debug/config.lua +++ b/nerv/examples/network_debug/config.lua @@ -12,7 +12,7 @@ function get_global_conf() layer_num = 1, chunk_size = 15, batch_size = 20, - max_iter = 1, + max_iter = 35, param_random = function() return (math.random() / 5 - 0.1) end, dropout_rate = 0.5, timer = nerv.Timer(), diff --git a/nerv/layer/duplicate.lua b/nerv/layer/duplicate.lua index 137472b..2621cdf 100644 --- a/nerv/layer/duplicate.lua +++ b/nerv/layer/duplicate.lua @@ -20,10 +20,7 @@ function DuplicateLayer:batch_resize() end function DuplicateLayer:propagate(input, output) - for i = 1, #self.dim_out do - output[i]:copy_from(input[1]) - -- FIXME: use reference copy to speed up - end + -- do nothing, use reference copy in nn/network.lua end function DuplicateLayer:back_propagate(bp_err, next_bp_err) diff --git a/nerv/nn/network.lua b/nerv/nn/network.lua index 910cdad..b06028e 100644 --- a/nerv/nn/network.lua +++ b/nerv/nn/network.lua @@ -228,9 +228,11 @@ function network:make_initial_store() err_memory[t][i][j] = self.mat_type(self.batch_size, dim_in[j]) err_memory[t][i][j]:fill(0) end - for j = 1, #dim_out do - memory[t][i][j] = self.mat_type(self.batch_size, dim_out[j]) - memory[t][i][j]:fill(self.nn_act_default) + if t < 1 or t > self.chunk_size or not nerv.is_type(self.layers[i], 'nerv.DuplicateLayer') then + for j = 1, #dim_out do + memory[t][i][j] = self.mat_type(self.batch_size, dim_out[j]) + memory[t][i][j]:fill(self.nn_act_default) + end end end if t < 1 or t > self.chunk_size then @@ -288,6 +290,28 @@ function network:make_initial_store() end end + -- reference copy for duplicate layer + for i = 1, #self.queue do + local t, id = self.queue[i].chunk, self.queue[i].id + if nerv.is_type(self.layers[id], 'nerv.DuplicateLayer') then + local _, dim_out = self.layers[id]:get_dim() + for j = 1, #dim_out do + if self.output[t][id][j] ~= nil then + nerv.error('duplicate output reference not nil') + end + self.output[t][id][j] = self.input[t][id][1] + local edge = self.output_conn[id][j] + local to, port, time = edge[1], edge[2], edge[3] + t + if time >= 1 and time <= self.chunk_size then + if self.input[time][to][port] ~= nil then + nerv.error('duplicate input reference not nil') + end + self.input[time][to][port] = self.output[t][id][j] + end + end + end + end + -- check dangling reference for t = 1, self.chunk_size do for i = 1, #self.dim_in do -- cgit v1.2.3 From 3dd235c8b6ea7ef275381866d11d1be828d27a06 Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Tue, 15 Mar 2016 12:48:07 +0800 Subject: fix duplicate bug on & --- nerv/examples/network_debug/config.lua | 10 +++++++-- nerv/examples/network_debug/reader.lua | 4 ++-- nerv/nn/network.lua | 40 ++++++++++++++++++++++++++++++++-- 3 files changed, 48 insertions(+), 6 deletions(-) diff --git a/nerv/examples/network_debug/config.lua b/nerv/examples/network_debug/config.lua index e20d5a9..0429e9a 100644 --- a/nerv/examples/network_debug/config.lua +++ b/nerv/examples/network_debug/config.lua @@ -35,6 +35,10 @@ function get_layers(global_conf) ['nerv.SoftmaxCELayer'] = { softmax = {dim_in = {global_conf.vocab_size, global_conf.vocab_size}, dim_out = {1}, compressed = true}, }, + ['nerv.DuplicateLayer'] = { + dup1 = {dim_in = {1}, dim_out = {1}}, + dup2 = {dim_in = {1}, dim_out = {1}}, + }, } for i = 1, global_conf.layer_num do layers['nerv.LSTMLayer']['lstm' .. i] = {dim_in = {global_conf.hidden_size}, dim_out = {global_conf.hidden_size}, pr = pr} @@ -45,12 +49,14 @@ end function get_connections(global_conf) local connections = { - {'[1]', 'select[1]', 0}, + {'[1]', 'dup1[1]', 0}, + {'dup1[1]', 'select[1]', 0}, {'select[1]', 'lstm1[1]', 0}, {'dropout' .. global_conf.layer_num .. '[1]', 'output[1]', 0}, {'output[1]', 'softmax[1]', 0}, {'[2]', 'softmax[2]', 0}, - {'softmax[1]', '[1]', 0}, + {'softmax[1]', 'dup2[1]', 0}, + {'dup2[1]', '[1]', 0}, } for i = 1, global_conf.layer_num do table.insert(connections, {'lstm' .. i .. '[1]', 'dropout' .. i .. '[1]', 0}) diff --git a/nerv/examples/network_debug/reader.lua b/nerv/examples/network_debug/reader.lua index 76a78cf..70c0c97 100644 --- a/nerv/examples/network_debug/reader.lua +++ b/nerv/examples/network_debug/reader.lua @@ -32,8 +32,8 @@ end function Reader:get_seq(input_file) local f = io.open(input_file, 'r') self.seq = {} - while true do - -- for i = 1, 26 do + -- while true do + for i = 1, 26 do local seq = f:read() if seq == nil then break diff --git a/nerv/nn/network.lua b/nerv/nn/network.lua index b06028e..6f7fe10 100644 --- a/nerv/nn/network.lua +++ b/nerv/nn/network.lua @@ -27,7 +27,17 @@ function network:__init(id, global_conf, network_conf) if self.input_conn[id][port] ~= nil then nerv.error('duplicate edge') end - self.input_conn[id][port] = {0, i, time} + if nerv.is_type(self.layers[id], 'nerv.DuplicateLayer') then + local tmp = nerv.IdentityLayer('', self.gconf, {dim_in = {self.dim_in[i]}, dim_out = {self.dim_in[i]}}) + table.insert(self.layers, tmp) + local new_id = #self.layers + self.input_conn[new_id] = {{0, i, time}} + self.output_conn[new_id] = {{id, port, 0}} + self.input_conn[id][port] = {new_id, 1, 0} + self.socket.inputs[i] = {new_id, 1, time} + else + self.input_conn[id][port] = {0, i, time} + end end for i = 1, #self.dim_out do local edge = self.socket.outputs[i] @@ -35,7 +45,17 @@ function network:__init(id, global_conf, network_conf) if self.output_conn[id][port] ~= nil then nerv.error('duplicate edge') end - self.output_conn[id][port] = {0, i, time} + if nerv.is_type(self.layers[id], 'nerv.DuplicateLayer') then + local tmp = nerv.IdentityLayer('', self.gconf, {dim_in = {self.dim_out[i]}, dim_out = {self.dim_out[i]}}) + table.insert(self.layers, tmp) + local new_id = #self.layers + self.input_conn[new_id] = {{id, port, 0}} + self.output_conn[new_id] = {{0, i, time}} + self.output_conn[id][port] = {new_id, 1, 0} + self.socket.outputs[i] = {new_id, 1, time} + else + self.output_conn[id][port] = {0, i, time} + end end self.delay = 0 @@ -140,8 +160,10 @@ function network:init(batch_size, chunk_size) collectgarbage('collect') self.flush = {} + self.gconf.mask = {} for t = 1, self.chunk_size do self.flush[t] = {} + self.gconf.mask[t] = self.mat_type(self.batch_size, 1) end end @@ -348,6 +370,7 @@ function network:make_initial_store() local dim_in, dim_out = self.layers[i]:get_dim() for j = 1, #dim_in do if self.input[t][i][j] == nil then + print(t,i,j,self.layers[i].id) nerv.error('input reference dangling') end if self.err_output[t][i][j] == nil then @@ -450,6 +473,19 @@ function network:mini_batch_init(info) self:set_err_output(self.info.err_output) end + -- calculate mask + for t = 1, self.chunk_size do + local tmp = self.gconf.mmat_type(self.batch_size, 1) + for i = 1, self.batch_size do + if t <= self.info.seq_length[i] then + tmp[i - 1][0] = 1 + else + tmp[i - 1][0] = 0 + end + end + self.gconf.mask[t]:copy_fromh(tmp) + end + -- calculate border self.max_length = 0 self.timestamp = self.timestamp + 1 -- cgit v1.2.3 From b547dd2a30e91ce124d50f763997070ea67c6f7e Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Tue, 15 Mar 2016 13:17:45 +0800 Subject: add mask on softmax_ce --- nerv/examples/network_debug/network.lua | 8 ++++---- nerv/layer/softmax_ce.lua | 4 +++- nerv/lib/matrix/generic/cukernel.cu | 22 ++++++++++++++++++++++ nerv/lib/matrix/generic/cumatrix.c | 14 +++++++++++++- nerv/lib/matrix/generic/cumatrix.h | 2 ++ nerv/lib/matrix/generic/mmatrix.c | 23 +++++++++++++++++++++++ nerv/lib/matrix/generic/mmatrix.h | 2 ++ nerv/matrix/generic/cumatrix.c | 1 + nerv/matrix/generic/matrix.c | 12 ++++++++++++ nerv/matrix/generic/mmatrix.c | 1 + 10 files changed, 83 insertions(+), 6 deletions(-) diff --git a/nerv/examples/network_debug/network.lua b/nerv/examples/network_debug/network.lua index 1841d21..386c3b0 100644 --- a/nerv/examples/network_debug/network.lua +++ b/nerv/examples/network_debug/network.lua @@ -72,12 +72,12 @@ function nn:process(data, do_train, reader) for t = 1, self.gconf.chunk_size do local tmp = info.output[t][1]:new_to_host() for i = 1, self.gconf.batch_size do - if t <= info.seq_length[i] then - total_err = total_err + math.log10(math.exp(tmp[i - 1][0])) - total_frame = total_frame + 1 - end + total_err = total_err + math.log10(math.exp(tmp[i - 1][0])) end end + for i = 1, self.gconf.batch_size do + total_frame = total_frame + info.seq_length[i] + end timer:toc('IO') timer:tic('network') diff --git a/nerv/layer/softmax_ce.lua b/nerv/layer/softmax_ce.lua index 7b4a80c..acd4ee6 100644 --- a/nerv/layer/softmax_ce.lua +++ b/nerv/layer/softmax_ce.lua @@ -61,14 +61,16 @@ function SoftmaxCELayer:propagate(input, output, t) end ce:mul_elem(ce, label) ce = ce:rowsum() + ce:set_values_by_mask(self.gconf.mask[t], 0) if output[1] ~= nil then output[1]:copy_from(ce) end -- add total ce self.total_ce = self.total_ce - ce:colsum()[0][0] - self.total_frames = self.total_frames + softmax:nrow() + self.total_frames = self.total_frames + self.gconf.mask[t]:colsum()[0][0] -- TODO: add colsame for uncompressed label if self.compressed then + classified:set_values_by_mask(self.gconf.mask[t], -1) self.total_correct = self.total_correct + classified:colsame(input[2])[0][0] end end diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu index 93121dc..4717209 100644 --- a/nerv/lib/matrix/generic/cukernel.cu +++ b/nerv/lib/matrix/generic/cukernel.cu @@ -324,6 +324,15 @@ __global__ void cudak_(rearrange_frm)(const MATRIX_ELEM *a, MATRIX_ELEM *b, b[j + i * stride] = a[j / step + (j % step) * orig_dim + i * stride]; } +__global__ void cudak_(set_values_by_mask)(const MATRIX_ELEM *a, MATRIX_ELEM *b, + int nrow, int ncol, + int astride, int bstride, double val) { + int j = blockIdx.x * blockDim.x + threadIdx.x; + int i = blockIdx.y * blockDim.y + threadIdx.y; + if (i >= nrow || j >= ncol || a[i * astride] != 0.0) return; + b[j + i * bstride] = val; +} + __global__ void cudak_(scale_rows_by_col)(const MATRIX_ELEM *a, MATRIX_ELEM *b, int nrow, int ncol, int astride, int bstride) { @@ -766,6 +775,19 @@ extern "C" { cudaStreamSynchronize(0); } + void cudak_(cuda_set_values_by_mask)(const Matrix *a, Matrix *b, double val) { + dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N); + dim3 numBlocks(CEIL_DIV(b->ncol, threadsPerBlock.x), + CEIL_DIV(b->nrow, threadsPerBlock.y)); + cudak_(set_values_by_mask)<<>> \ + (MATRIX_ELEM_PTR(a), MATRIX_ELEM_PTR(b), + b->nrow, b->ncol, + a->stride / sizeof(MATRIX_ELEM), + b->stride / sizeof(MATRIX_ELEM), + val); + cudaStreamSynchronize(0); + } + void cudak_(cuda_scale_rows_by_row)(const Matrix *a, Matrix *b) { dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N); dim3 numBlocks(CEIL_DIV(b->ncol, threadsPerBlock.x), diff --git a/nerv/lib/matrix/generic/cumatrix.c b/nerv/lib/matrix/generic/cumatrix.c index 6d84663..bc5f285 100644 --- a/nerv/lib/matrix/generic/cumatrix.c +++ b/nerv/lib/matrix/generic/cumatrix.c @@ -515,7 +515,7 @@ void nerv_matrix_(prefixsum_row)(Matrix *a, const Matrix *b, NERV_SET_STATUS(status, NERV_NORMAL, 0); } -void nerv_matrix_(diagonalize)(Matrix *a, CuContext * context, Status *status) { +void nerv_matrix_(diagonalize)(Matrix *a, CuContext *context, Status *status) { if (a->nrow != a->ncol) NERV_EXIT_STATUS(status, MAT_MISMATCH_DIM, 0); PROFILE_START @@ -524,6 +524,18 @@ void nerv_matrix_(diagonalize)(Matrix *a, CuContext * context, Status *status) { NERV_SET_STATUS(status, NERV_NORMAL, 0); } +void nerv_matrix_(set_values_by_mask)(Matrix *a, const Matrix *b, double val, + CuContext *context, Status *status) { + if (a->nrow != b->nrow) + NERV_EXIT_STATUS(status, MAT_MISMATCH_DIM, 0); + if (b->ncol != 1) + NERV_EXIT_STATUS(status, MAT_COL_VECTOR_EXP, 0); + PROFILE_START + cudak_(cuda_set_values_by_mask)(b, a, val); + PROFILE_STOP + NERV_SET_STATUS(status, NERV_NORMAL, 0); +} + static void cuda_matrix_(free)(MATRIX_ELEM *ptr, CuContext *context, Status *status) { CUDA_SAFE_SYNC_CALL(cudaFree(ptr), status); NERV_SET_STATUS(status, NERV_NORMAL, 0); diff --git a/nerv/lib/matrix/generic/cumatrix.h b/nerv/lib/matrix/generic/cumatrix.h index de3a09e..79bfc76 100644 --- a/nerv/lib/matrix/generic/cumatrix.h +++ b/nerv/lib/matrix/generic/cumatrix.h @@ -35,6 +35,8 @@ void nerv_matrix_(fill)(Matrix *self, double val, CuContext *context, Status *status); void nerv_matrix_(diagonalize)(Matrix *self, CuContext *context, Status *status); +void nerv_matrix_(set_values_by_mask)(Matrix *self, Matrix *mask, double val, + CuContext *context, Status *status); void nerv_matrix_(copy_fromd)(Matrix *a, const Matrix *b, int a_begin, int b_begin, int b_end, CuContext *context, Status *status); diff --git a/nerv/lib/matrix/generic/mmatrix.c b/nerv/lib/matrix/generic/mmatrix.c index badddbd..e356de7 100644 --- a/nerv/lib/matrix/generic/mmatrix.c +++ b/nerv/lib/matrix/generic/mmatrix.c @@ -507,6 +507,29 @@ void nerv_matrix_(scale_rows_by_col)(Matrix *a, const Matrix *b, NERV_SET_STATUS(status, NERV_NORMAL, 0); } +void nerv_matrix_(set_values_by_mask)(Matrix *a, const Matrix *b, double val, + MContext *context, Status *status) { + if (a->nrow != b->nrow) + NERV_EXIT_STATUS(status, MAT_MISMATCH_DIM, 0); + if (b->ncol != 1) + NERV_EXIT_STATUS(status, MAT_COL_VECTOR_EXP, 0); + int i, j; + size_t astride = a->stride, bstride = b->stride; + MATRIX_ELEM *arow = MATRIX_ELEM_PTR(a), + *brow = MATRIX_ELEM_PTR(b); + for (i = 0; i < a->nrow; i++) + { + if (brow[0] == 0.0) + { + for (j = 0; j < a->ncol; j++) + arow[j] = val; + } + arow = MATRIX_NEXT_ROW_PTR(arow, astride); + brow = MATRIX_NEXT_ROW_PTR(brow, bstride); + } + NERV_SET_STATUS(status, NERV_NORMAL, 0); +} + static void host_matrix_(free)(MATRIX_ELEM *ptr, MContext *context, Status *status) { free(ptr); NERV_SET_STATUS(status, NERV_NORMAL, 0); diff --git a/nerv/lib/matrix/generic/mmatrix.h b/nerv/lib/matrix/generic/mmatrix.h index 6d17c99..41c39f6 100644 --- a/nerv/lib/matrix/generic/mmatrix.h +++ b/nerv/lib/matrix/generic/mmatrix.h @@ -48,6 +48,8 @@ void nerv_matrix_(expand_frm)(Matrix *a, const Matrix *b, int cont, MContext *context, Status *status); void nerv_matrix_(rearrange_frm)(Matrix *a, const Matrix *b, int step, MContext *context, Status *status); +void nerv_matrix_(set_values_by_mask)(Matrix *a, const Matrix *b, double val, + MContext *context, Status *status); void nerv_matrix_(scale_rows_by_col)(Matrix *a, const Matrix *b, MContext *context, Status *status); void nerv_matrix_(scale_rows_by_row)(Matrix *a, const Matrix *b, diff --git a/nerv/matrix/generic/cumatrix.c b/nerv/matrix/generic/cumatrix.c index 0c90d39..9577fd5 100644 --- a/nerv/matrix/generic/cumatrix.c +++ b/nerv/matrix/generic/cumatrix.c @@ -268,6 +268,7 @@ static const luaL_Reg nerv_matrix_(extra_methods)[] = { {"scale_rows_by_col", nerv_matrix_(lua_scale_rows_by_col)}, {"prefixsum_row", nerv_matrix_(lua_prefixsum_row)}, {"diagonalize", nerv_matrix_(lua_diagonalize)}, + {"set_values_by_mask", nerv_matrix_(lua_set_values_by_mask)}, #ifdef __NERV_FUTURE_CUDA_7 {"update_select_rows_by_rowidx", nerv_matrix_(lua_update_select_rows_by_rowidx)}, {"update_select_rows_by_colidx", nerv_matrix_(lua_update_select_rows_by_colidx)}, diff --git a/nerv/matrix/generic/matrix.c b/nerv/matrix/generic/matrix.c index fe07585..3e91933 100644 --- a/nerv/matrix/generic/matrix.c +++ b/nerv/matrix/generic/matrix.c @@ -395,4 +395,16 @@ static int nerv_matrix_(lua_diagonalize)(lua_State *L) { return 0; } +static int nerv_matrix_(lua_set_values_by_mask)(lua_State *L) { + Status status; + MATRIX_CONTEXT *context; + MATRIX_GET_CONTEXT(L, 4); + Matrix *a = luaT_checkudata(L, 1, nerv_matrix_(tname)); + Matrix *mask = luaT_checkudata(L, 2, nerv_matrix_(tname)); + double val = luaL_checknumber(L, 3); + nerv_matrix_(set_values_by_mask)(a, mask, val, context, &status); + NERV_LUA_CHECK_STATUS(L, status); + return 0; +} + #endif diff --git a/nerv/matrix/generic/mmatrix.c b/nerv/matrix/generic/mmatrix.c index a5e5969..de1eaa3 100644 --- a/nerv/matrix/generic/mmatrix.c +++ b/nerv/matrix/generic/mmatrix.c @@ -117,6 +117,7 @@ static const luaL_Reg nerv_matrix_(extra_methods)[] = { {"clip", nerv_matrix_(lua_clip)}, {"fill", nerv_matrix_(lua_fill)}, {"diagonalize", nerv_matrix_(lua_diagonalize)}, + {"set_values_by_mask", nerv_matrix_(lua_set_values_by_mask)}, {"sigmoid", nerv_matrix_(lua_sigmoid)}, {"sigmoid_grad", nerv_matrix_(lua_sigmoid_grad)}, {"softmax", nerv_matrix_(lua_softmax)}, -- cgit v1.2.3 From e15307f071813e2eb56f7f83229b91141961325a Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Tue, 15 Mar 2016 13:30:33 +0800 Subject: add total time on profile --- nerv/examples/network_debug/reader.lua | 4 ++-- nerv/lib/matrix/cumatrix.c | 6 +++++- nerv/lib/matrix/mmatrix.c | 6 +++++- 3 files changed, 12 insertions(+), 4 deletions(-) diff --git a/nerv/examples/network_debug/reader.lua b/nerv/examples/network_debug/reader.lua index 70c0c97..76a78cf 100644 --- a/nerv/examples/network_debug/reader.lua +++ b/nerv/examples/network_debug/reader.lua @@ -32,8 +32,8 @@ end function Reader:get_seq(input_file) local f = io.open(input_file, 'r') self.seq = {} - -- while true do - for i = 1, 26 do + while true do + -- for i = 1, 26 do local seq = f:read() if seq == nil then break diff --git a/nerv/lib/matrix/cumatrix.c b/nerv/lib/matrix/cumatrix.c index aec4d60..43448bf 100644 --- a/nerv/lib/matrix/cumatrix.c +++ b/nerv/lib/matrix/cumatrix.c @@ -9,15 +9,19 @@ void nerv_cuda_context_print_profile(CuContext *context) { HashMap *profile = context->profile; size_t i; + float tmp, tot = 0; fprintf(stderr, "*** [nerv cumatrix profile] **\n"); for (i = 0; i < profile->size; i++) { HashNode *ptr; for (ptr = profile->bucket[i]; ptr; ptr = ptr->next) { - fprintf(stderr, "%s:\t%.6f\n", ptr->key, *(float *)ptr->val); + tmp = *(float *)ptr->val; + fprintf(stderr, "%s:\t%.6f\n", ptr->key, tmp); + tot += tmp; } } + fprintf(stderr, "Total time:\t%.6f\n", tot); } void nerv_cuda_context_clear_profile(CuContext *context) { diff --git a/nerv/lib/matrix/mmatrix.c b/nerv/lib/matrix/mmatrix.c index e40b160..006735d 100644 --- a/nerv/lib/matrix/mmatrix.c +++ b/nerv/lib/matrix/mmatrix.c @@ -8,15 +8,19 @@ void nerv_host_context_print_profile(MContext *context) { HashMap *profile = context->profile; size_t i; + float tmp, tot = 0; fprintf(stderr, "*** [nerv mmatrix profile] **\n"); for (i = 0; i < profile->size; i++) { HashNode *ptr; for (ptr = profile->bucket[i]; ptr; ptr = ptr->next) { - fprintf(stderr, "%s:\t%.6f\n", ptr->key, *(float *)ptr->val); + tmp = *(float *)ptr->val; + fprintf(stderr, "%s:\t%.6f\n", ptr->key, tmp); + tot += tmp; } } + fprintf(stderr, "Total time:\t%.6f\n", tot); } void nerv_host_context_clear_profile(MContext *context) { -- cgit v1.2.3