Merge pull request #37 from cloudygoose/master

add matrix:clip & affine_recurrent layer

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)},