From 72acf24e248cca7d69658d02939f99d57d02e9a9 Mon Sep 17 00:00:00 2001 From: cloudygoose Date: Wed, 8 Jul 2015 08:54:01 +0800 Subject: added matrix:clip, TODO:testing --- nerv/lib/matrix/cukernel.h | 1 + nerv/lib/matrix/generic/cukernel.cu | 22 ++++++++++++++++++++++ nerv/lib/matrix/generic/cumatrix.c | 7 +++++++ nerv/lib/matrix/generic/cumatrix.h | 1 + nerv/matrix/generic/cukernel.cu | 21 +++++++++++++++++++++ nerv/matrix/generic/cumatrix.c | 11 +++++++++++ 6 files changed, 63 insertions(+) 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)<<>> \ + (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)<<>> \ + (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)}, -- cgit v1.2.3-70-g09d2