aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--.gitignore1
-rw-r--r--nerv/Makefile8
-rw-r--r--nerv/layer/affine_recurrent.lua87
-rw-r--r--nerv/layer/init.lua1
-rw-r--r--nerv/lib/matrix/cukernel.h1
-rw-r--r--nerv/lib/matrix/generic/cukernel.cu22
-rw-r--r--nerv/lib/matrix/generic/cumatrix.c7
-rw-r--r--nerv/lib/matrix/generic/cumatrix.h1
-rw-r--r--nerv/matrix/generic/cukernel.cu21
-rw-r--r--nerv/matrix/generic/cumatrix.c11
10 files changed, 156 insertions, 4 deletions
diff --git a/.gitignore b/.gitignore
index 24f2f11..5c87b71 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,4 +1,5 @@
*.o
install/
+build/
*.swp
*.swo
diff --git a/nerv/Makefile b/nerv/Makefile
index 224cc8a..022e2fb 100644
--- a/nerv/Makefile
+++ b/nerv/Makefile
@@ -31,13 +31,13 @@ OBJS := $(CORE_OBJS) $(NERV_OBJS) $(LUAT_OBJS)
LIBS := $(INST_LIBDIR)/libnerv.so $(LIB_PATH)/libnervcore.so $(LIB_PATH)/libluaT.so
LUA_LIBS := matrix/init.lua io/init.lua init.lua \
layer/init.lua layer/affine.lua layer/sigmoid.lua layer/softmax_ce.lua \
- layer/window.lua layer/bias.lua layer/combiner.lua layer/mse.lua \
+ layer/window.lua layer/bias.lua layer/combiner.lua layer/mse.lua layer/affine_recurrent.lua\
nn/init.lua nn/layer_repo.lua nn/param_repo.lua nn/layer_dag.lua \
- io/sgd_buffer.lua
+ io/sgd_buffer.lua
INCLUDE := -I $(LUA_INCDIR) -DLUA_USE_APICHECK
-#CUDA_BASE := /usr/local/cuda-6.5
-CUDA_BASE := /usr/local/cuda-5.0
+CUDA_BASE := /usr/local/cuda-6.5
+#CUDA_BASE := /usr/local/cuda-5.0
CUDA_INCLUDE := -I $(CUDA_BASE)/include/
INCLUDE += $(CUDA_INCLUDE)
diff --git a/nerv/layer/affine_recurrent.lua b/nerv/layer/affine_recurrent.lua
new file mode 100644
index 0000000..59d259c
--- /dev/null
+++ b/nerv/layer/affine_recurrent.lua
@@ -0,0 +1,87 @@
+local Recurrent = nerv.class('nerv.AffineRecurrentLayer', 'nerv.Layer')
+
+--id: string
+--global_conf: table
+--layer_conf: table
+--Get Parameters
+function Recurrent:__init(id, global_conf, layer_conf)
+ self.id = id
+ self.dim_in = layer_conf.dim_in
+ self.dim_out = layer_conf.dim_out
+ self.gconf = global_conf
+
+ self.bp = layer_conf.bp
+ self.ltp_hh = layer_conf.ltp_hh --from hidden to hidden
+
+ self:check_dim_len(2, 1)
+ self.direct_update = layer_conf.direct_update
+
+ self.clip = layer_conf.clip --clip error in back_propagate
+end
+
+--Check parameter
+function Recurrent:init(batch_size)
+ if (self.ltp_hh.trans:ncol() ~= self.bp.trans:ncol()) then
+ nerv.error("mismatching dimensions of ltp and bp")
+ end
+ if (self.dim_in[1] ~= self.ltp_hh.trans:nrow() or
+ self.dim_in[2] ~= self.ltp_hh.trans:nrow()) then
+ nerv.error("mismatching dimensions of ltp and input")
+ end
+ if (self.dim_out[1] ~= self.bp.trans:ncol()) then
+ nerv.error("mismatching dimensions of bp and output")
+ end
+
+ self.ltp_hh_grad = self.ltp_hh.trans:create()
+ self.ltp_hh:train_init()
+ self.bp:train_init()
+end
+
+function Recurrent:update(bp_err, input, output)
+ if (self.direct_update == true) then
+ local ltp_hh = self.ltp_hh.trans
+ local bp = self.bp.trans
+ local gconf = self.gconf
+ -- momentum gain
+ local mmt_gain = 1.0 / (1.0 - gconf.momentum);
+ local n = input[1]:nrow() * mmt_gain
+ -- update corrections (accumulated errors)
+ self.ltp_hh.correction:mul(input[2], bp_err[1], 1.0, gconf.momentum, 'T', 'N')
+ self.bp.correction:add(bc, bp_err[1]:colsum(), gconf.momentum, 1.0)
+ -- perform update
+ ltp_hh:add(ltp_hh, self.ltp_hh.correction, 1.0, -gconf.lrate / n)
+ bp:add(bp, self.bp.correction, 1.0, -gconf.lrate / n)
+ -- weight decay
+ ltp_hh:add(ltp_hh, ltp_hh, 1.0, -gconf.lrate * gconf.wcost)
+ else
+ self.ltp_hh_grad:mul(input[2], bp_err[1], 1.0, 0.0, 'T', 'N')
+ self.ltp_hh:update(self.ltp_hh_grad)
+ self.bp:update(bp_err[1]:colsum())
+ end
+end
+
+function Recurrent:propagate(input, output)
+ output[1]:copy_fromd(input[1])
+ output[1]:mul(input[2], self.ltp_hh.trans, 1.0, 1.0, 'N', 'N')
+ output[1]:add_row(self.bp.trans, 1.0)
+end
+
+function Recurrent:back_propagate(bp_err, next_bp_err, input, output)
+ next_bp_err[1]:copy_fromd(bp_err[1])
+ next_bp_err[2]:mul(bp_err[1], self.ltp_hh.trans, 1.0, 0.0, 'N', 'T')
+ --[[
+ for i = 0, next_bp_err[2]:nrow() - 1 do
+ for j = 0, next_bp_err[2]:ncol() - 1 do
+ if (next_bp_err[2][i][j] > 10) then next_bp_err[2][i][j] = 10 end
+ if (next_bp_err[2][i][j] < -10) then next_bp_err[2][i][j] = -10 end
+ end
+ end
+ ]]--
+ if (self.clip ~= nil) then
+ next_bp_err[2]:clip(-self.clip, self.clip)
+ end
+end
+
+function Recurrent:get_params()
+ return nerv.ParamRepo({self.ltp_hh, self.bp})
+end
diff --git a/nerv/layer/init.lua b/nerv/layer/init.lua
index 3c55a94..7172f99 100644
--- a/nerv/layer/init.lua
+++ b/nerv/layer/init.lua
@@ -77,3 +77,4 @@ nerv.include('bias.lua')
nerv.include('window.lua')
nerv.include('mse.lua')
nerv.include('combiner.lua')
+nerv.include('affine_recurrent.lua')
diff --git a/nerv/lib/matrix/cukernel.h b/nerv/lib/matrix/cukernel.h
index 8a1494f..7bb4c2c 100644
--- a/nerv/lib/matrix/cukernel.h
+++ b/nerv/lib/matrix/cukernel.h
@@ -12,6 +12,7 @@ void cudak_(cuda_softmax_denominator)(const Matrix *a, const Matrix *max, Matrix
void cudak_(cuda_softmax_final)(const Matrix *a, const Matrix *max, const Matrix *deno, Matrix *b);
void cudak_(cuda_add_row)(const Matrix *a, Matrix *b, double beta);
void cudak_(cuda_fill)(Matrix *a, double val);
+void cudak_(cuda_clip)(Matrix *a, double val_1, double val_2);
void cudak_(cuda_expand_frm)(const Matrix *a, Matrix *b, int context);
void cudak_(cuda_rearrange_frm)(const Matrix *a, Matrix *b, int step);
void cudak_(cuda_scale_rows_by_row)(const Matrix *a, Matrix *b);
diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu
index 6111193..e337798 100644
--- a/nerv/lib/matrix/generic/cukernel.cu
+++ b/nerv/lib/matrix/generic/cukernel.cu
@@ -213,6 +213,18 @@ __global__ void cudak_(fill)(MATRIX_ELEM *a,
a[j + i * stride] = val;
}
+__global__ void cudak_(clip)(MATRIX_ELEM *a,
+ int nrow, int ncol, int stride, double val_1, double val_2) {
+ int j = blockIdx.x * blockDim.x + threadIdx.x;
+ int i = blockIdx.y * blockDim.y + threadIdx.y;
+ if (i >= nrow || j >= ncol) return;
+ if (a[j + i * stride] > val_2)
+ a[j + i * stride] = val_2;
+ else
+ if (a[j + i * stride] < val_1)
+ a[j + i * stride] = val_1;
+}
+
__global__ void cudak_(expand_frm)(const MATRIX_ELEM *a, MATRIX_ELEM *b,
int nrow, int ncol,
int enrow, int encol,
@@ -510,6 +522,16 @@ extern "C" {
cudaStreamSynchronize(0);
}
+ void cudak_(cuda_clip)(Matrix *a, double val_1, double val_2) {
+ dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N);
+ dim3 numBlocks(CEIL_DIV(a->ncol, threadsPerBlock.x),
+ CEIL_DIV(a->nrow, threadsPerBlock.y));
+ cudak_(clip)<<<numBlocks, threadsPerBlock>>> \
+ (MATRIX_ELEM_PTR(a), a->nrow, a->ncol,
+ a->stride / sizeof(MATRIX_ELEM), val_1, val_2);
+ cudaStreamSynchronize(0);
+ }
+
void cudak_(cuda_expand_frm)(const Matrix *a, Matrix *b, int context) {
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 772b78d..40a0030 100644
--- a/nerv/lib/matrix/generic/cumatrix.c
+++ b/nerv/lib/matrix/generic/cumatrix.c
@@ -189,6 +189,13 @@ void nerv_matrix_(fill)(Matrix *self, double val, Status *status) {
NERV_SET_STATUS(status, NERV_NORMAL, 0);
}
+void nerv_matrix_(clip)(Matrix *self, double val_1, double val_2, Status *status) {
+ PROFILE_START
+ cudak_(cuda_clip)(self, val_1, val_2);
+ PROFILE_STOP
+ NERV_SET_STATUS(status, NERV_NORMAL, 0);
+}
+
void nerv_matrix_(copy_fromd)(Matrix *a, const Matrix *b,
int a_begin, int b_begin, int b_end,
Status *status) {
diff --git a/nerv/lib/matrix/generic/cumatrix.h b/nerv/lib/matrix/generic/cumatrix.h
index 5cfe9d5..3f1f8a3 100644
--- a/nerv/lib/matrix/generic/cumatrix.h
+++ b/nerv/lib/matrix/generic/cumatrix.h
@@ -20,6 +20,7 @@ void nerv_matrix_(rowmax_idx)(Matrix *a, Matrix **b, Matrix **idx,
Status *status);
void nerv_matrix_(add_row)(Matrix *b, const Matrix *a, double beta,
Status *status);
+void nerv_matrix_(clip)(Matrix *self, double val_1, double val_2, Status *status);
void nerv_matrix_(fill)(Matrix *self, double val, Status *status);
void nerv_matrix_(copy_fromd)(Matrix *a, const Matrix *b,
int a_begin, int b_begin, int b_end,
diff --git a/nerv/matrix/generic/cukernel.cu b/nerv/matrix/generic/cukernel.cu
index d6c8adc..2ae5e62 100644
--- a/nerv/matrix/generic/cukernel.cu
+++ b/nerv/matrix/generic/cukernel.cu
@@ -213,6 +213,17 @@ __global__ void cudak_(fill)(MATRIX_ELEM *a,
a[j + i * stride] = val;
}
+__global__ void cudak_(clip)(MATRIX_ELEM *a,
+ int nrow, int ncol, int stride, double val_1, double val_2) {
+ int j = blockIdx.x * blockDim.x + threadIdx.x;
+ int i = blockIdx.y * blockDim.y + threadIdx.y;
+ if (i >= nrow || j >= ncol) return;
+ if (a[j + i * stride] > val_2)
+ a[j + i * stride] = val_2;
+ else if (a[j + i * stride] < val_1)
+ a[j + i * stride] = val_1;
+}
+
__global__ void cudak_(expand_frm)(const MATRIX_ELEM *a, MATRIX_ELEM *b,
int nrow, int ncol,
int enrow, int encol,
@@ -510,6 +521,16 @@ extern "C" {
cudaStreamSynchronize(0);
}
+ void cudak_(cuda_clip)(Matrix *a, double val_1, double val_2) {
+ dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N);
+ dim3 numBlocks(CEIL_DIV(a->ncol, threadsPerBlock.x),
+ CEIL_DIV(a->nrow, threadsPerBlock.y));
+ cudak_(clip)<<<numBlocks, threadsPerBlock>>> \
+ (MATRIX_ELEM_PTR(a), a->nrow, a->ncol,
+ a->stride / sizeof(MATRIX_ELEM), val_1, val_2);
+ cudaStreamSynchronize(0);
+ }
+
void cudak_(cuda_expand_frm)(const Matrix *a, Matrix *b, int context) {
dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N);
dim3 numBlocks(CEIL_DIV(b->ncol, threadsPerBlock.x),
diff --git a/nerv/matrix/generic/cumatrix.c b/nerv/matrix/generic/cumatrix.c
index 311b503..4bdf5f0 100644
--- a/nerv/matrix/generic/cumatrix.c
+++ b/nerv/matrix/generic/cumatrix.c
@@ -149,6 +149,16 @@ static int nerv_matrix_(lua_fill)(lua_State *L) {
return 0;
}
+static int nerv_matrix_(lua_clip)(lua_State *L) {
+ Status status;
+ Matrix *self = luaT_checkudata(L, 1, nerv_matrix_(tname));
+ double val_1 = luaL_checknumber(L, 2);
+ double val_2 = luaL_checknumber(L, 3);
+ nerv_matrix_(clip)(self, val_1, val_2, &status);
+ NERV_LUA_CHECK_STATUS(L, status);
+ return 0;
+}
+
static int nerv_matrix_(lua_copy_fromd)(lua_State *L) {
Status status;
Matrix *a = luaT_checkudata(L, 1, nerv_matrix_(tname));
@@ -294,6 +304,7 @@ static const luaL_Reg nerv_matrix_(extra_methods)[] = {
{"add", nerv_matrix_(lua_add)},
{"mul", nerv_matrix_(lua_mul)},
{"add_row", nerv_matrix_(lua_add_row)},
+ {"clip", nerv_matrix_(lua_clip)},
{"fill", nerv_matrix_(lua_fill)},
{"sigmoid", nerv_matrix_(lua_sigmoid)},
{"sigmoid_grad", nerv_matrix_(lua_sigmoid_grad)},
Powered by cgit, Themed by Ted Yin.