From a51498d2761714f4e034f036b84b4b89a88e9539 Mon Sep 17 00:00:00 2001 From: Qi Liu Date: Wed, 24 Feb 2016 16:16:03 +0800 Subject: update LSTM layer --- nerv/lib/matrix/generic/cukernel.cu | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) (limited to 'nerv/lib/matrix/generic/cukernel.cu') diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu index 51e3b6a..311a6ce 100644 --- a/nerv/lib/matrix/generic/cukernel.cu +++ b/nerv/lib/matrix/generic/cukernel.cu @@ -250,6 +250,14 @@ __global__ void cudak_(fill)(MATRIX_ELEM *a, a[j + i * stride] = val; } +__global__ void cudak_(diagonalize)(MATRIX_ELEM *a, + int nrow, int ncol, int stride) { + int j = blockIdx.x * blockDim.x + threadIdx.x; + int i = blockIdx.y * blockDim.y + threadIdx.y; + if (i >= nrow || j >= ncol || i == j) return; + a[j + i * stride] = 0; +} + __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; @@ -679,6 +687,16 @@ extern "C" { cudaStreamSynchronize(0); } + void cudak_(cuda_diagonalize)(Matrix *a) { + dim3 threadsPerBlock(CUDA_THREADS_N, CUDA_THREADS_N); + dim3 numBlocks(CEIL_DIV(a->ncol, threadsPerBlock.x), + CEIL_DIV(a->nrow, threadsPerBlock.y)); + cudak_(diagonalize)<<>> \ + (MATRIX_ELEM_PTR(a), a->nrow, a->ncol, + a->stride / sizeof(MATRIX_ELEM)); + 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), -- cgit v1.2.3