diff options
author | Ted Yin <Determinant@users.noreply.github.com> | 2016-03-12 13:17:38 +0800 |
---|---|---|
committer | Ted Yin <Determinant@users.noreply.github.com> | 2016-03-12 13:17:38 +0800 |
commit | 2b03555ea53a47e87d03a79feb866c868d424f01 (patch) | |
tree | 63cd01ee70d056d3a159a1e7d9aa4ea6f414d213 /nerv/lib/matrix/generic/cukernel.cu | |
parent | e8b1007d99691c08dd1b71f5733eb3cd2827dc64 (diff) | |
parent | 2660af7f6a9ac243a8ad38bf3375ef0fd292bf52 (diff) |
Merge pull request #31 from liuq901/master
modfiy param generate & rewrite LSTM layer
Diffstat (limited to 'nerv/lib/matrix/generic/cukernel.cu')
-rw-r--r-- | nerv/lib/matrix/generic/cukernel.cu | 18 |
1 files changed, 18 insertions, 0 deletions
diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu index 0e09cfa..93121dc 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; @@ -678,6 +686,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)<<<numBlocks, threadsPerBlock>>> \ + (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), |