From 20e0c009ab387107ed1a492dd42b0253ca19ed37 Mon Sep 17 00:00:00 2001 From: txh18 Date: Wed, 20 Jan 2016 14:26:54 +0800 Subject: the prefixsum_row operation of matrix is implemented --- nerv/lib/matrix/generic/cukernel.cu | 50 +++++++++++++++++++++++++++++++++++++ nerv/lib/matrix/generic/cumatrix.c | 8 ++++++ nerv/lib/matrix/generic/cumatrix.h | 1 + nerv/matrix/generic/cumatrix.c | 10 ++++++++ 4 files changed, 69 insertions(+) diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu index 2b696d5..995059c 100644 --- a/nerv/lib/matrix/generic/cukernel.cu +++ b/nerv/lib/matrix/generic/cukernel.cu @@ -376,6 +376,22 @@ __global__ void cudak_(copy_rows_by_colidx)(const MATRIX_ELEM *a, MATRIX_ELEM *b b[j + i * stride] = a[j + k * stride]; } +__global__ void cudak_(prefixsum_row_reduce)(const MATRIX_ELEM *a, MATRIX_ELEM *b, + int nrow, int ncol, int stride_a, int stride_b, int offset) { + int j = blockIdx.x * blockDim.x + threadIdx.x; + int i = blockIdx.y * blockDim.y + threadIdx.y; + long idx_a, idx_b; + if (i >= nrow || j >= ncol) return; + idx_b = j + i * stride_b; + idx_a = j + i * stride_a; + //b[idx] = 1.0 / (1.0 + exp(-a[idx])); + if (j >= offset) + b[idx_b] = a[idx_a] + a[idx_a - offset]; + else + b[idx_b] = a[idx_a]; +} + + extern "C" { #include "../cukernel.h" @@ -737,6 +753,40 @@ extern "C" { cudaStreamSynchronize(0); } + void cudak_(cuda_prefixsum_row)(const Matrix *a, Matrix *b) { + dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N); + dim3 numBlocks(CEIL_DIV(b->ncol, threadsPerBlock.x), + CEIL_DIV(b->nrow, threadsPerBlock.y)); + + MATRIX_ELEM *tmp[2]; + size_t tmp_stride[2]; + cudaMallocPitch(tmp, tmp_stride + 0, a->ncol * sizeof(MATRIX_ELEM), a->nrow); + cudaMallocPitch(tmp + 1, tmp_stride + 1, a->ncol * sizeof(MATRIX_ELEM), a->nrow); + + int offset = 1; + cudak_(prefixsum_row_reduce)<<>> \ + (MATRIX_ELEM_PTR(a), tmp[0], b->nrow, b->ncol, + a->stride / sizeof(MATRIX_ELEM), tmp_stride[0] / sizeof(MATRIX_ELEM), offset); + int pin = 0, pout = 1; + + for (offset = 2;offset <= a->ncol / 2;offset *= 2) { + cudak_(prefixsum_row_reduce)<<>> \ + (tmp[pin], tmp[pout], b->nrow, b->ncol, + tmp_stride[pin] / sizeof(MATRIX_ELEM), tmp_stride[pout] / sizeof(MATRIX_ELEM), offset); + pin = 1 - pin; + pout = 1 - pout; + } + + cudak_(prefixsum_row_reduce)<<>> \ + (tmp[pin], MATRIX_ELEM_PTR(b), b->nrow, b->ncol, + tmp_stride[pin] / sizeof(MATRIX_ELEM), b->stride / sizeof(MATRIX_ELEM), offset); + + cudaFree(tmp[0]); + cudaFree(tmp[1]); + + cudaStreamSynchronize(0); + } + void cudak_(cuda_decompress)(const Matrix *a, Matrix *b) { dim3 threadsPerBlock(1, CUDA_THREADS_NN); dim3 numBlocks(1, CEIL_DIV(a->nrow, threadsPerBlock.y)); diff --git a/nerv/lib/matrix/generic/cumatrix.c b/nerv/lib/matrix/generic/cumatrix.c index 7582725..dafadb2 100644 --- a/nerv/lib/matrix/generic/cumatrix.c +++ b/nerv/lib/matrix/generic/cumatrix.c @@ -486,6 +486,14 @@ void nerv_matrix_(scale_rows_by_row)(Matrix *a, const Matrix *b, NERV_SET_STATUS(status, NERV_NORMAL, 0); } +void nerv_matrix_(prefixsum_row)(Matrix *a, const Matrix *b, Status *status) { + CHECK_SAME_DIMENSION(a, b, status); + PROFILE_START + cudak_(cuda_prefixsum_row)(b, a); + PROFILE_STOP + NERV_SET_STATUS(status, NERV_NORMAL, 0); +} + static void cuda_matrix_(free)(MATRIX_ELEM *ptr, 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 e82dccd..41e9e35 100644 --- a/nerv/lib/matrix/generic/cumatrix.h +++ b/nerv/lib/matrix/generic/cumatrix.h @@ -61,3 +61,4 @@ void nerv_matrix_(scale_rows_by_col)(Matrix *a, const Matrix *b, Status *status); void nerv_matrix_(scale_rows_by_row)(Matrix *a, const Matrix *b, Status *status); +void nerv_matrix_(prefixsum_row)(Matrix *a, const Matrix *b, Status *status); diff --git a/nerv/matrix/generic/cumatrix.c b/nerv/matrix/generic/cumatrix.c index edd7b0a..29fa7b1 100644 --- a/nerv/matrix/generic/cumatrix.c +++ b/nerv/matrix/generic/cumatrix.c @@ -52,6 +52,15 @@ static int nerv_matrix_(lua_sigmoid)(lua_State *L) { return 0; } +static int nerv_matrix_(lua_prefixsum_row)(lua_State *L) { + Status status; + Matrix *a = luaT_checkudata(L, 1, nerv_matrix_(tname)); + Matrix *b = luaT_checkudata(L, 2, nerv_matrix_(tname)); + nerv_matrix_(prefixsum_row)(a, b, &status); + NERV_LUA_CHECK_STATUS(L, status); + return 0; +} + static int nerv_matrix_(lua_sigmoid_grad)(lua_State *L) { Status status; Matrix *nerr = luaT_checkudata(L, 1, nerv_matrix_(tname)); @@ -407,6 +416,7 @@ static const luaL_Reg nerv_matrix_(extra_methods)[] = { {"rearrange_frm", nerv_matrix_(lua_rearrange_frm)}, {"scale_rows_by_row", nerv_matrix_(lua_scale_rows_by_row)}, {"scale_rows_by_col", nerv_matrix_(lua_scale_rows_by_col)}, + {"prefixsum_row", nerv_matrix_(lua_prefixsum_row)}, #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)}, -- cgit v1.2.3