aboutsummaryrefslogtreecommitdiff
path: root/matrix/generic/cukernel.cu
diff options
context:
space:
mode:
authorDeterminant <[email protected]>2015-05-26 18:56:29 +0800
committerDeterminant <[email protected]>2015-05-26 18:56:29 +0800
commitc7e170d4f3cc9f73380380a96318cbea437d48ba (patch)
tree41abed014da976a4d27a1cbdcfa3171d45657eac /matrix/generic/cukernel.cu
parentfd389b72623dcb44009076c3819a74a79b6f94be (diff)
add fill for cumatrix
Diffstat (limited to 'matrix/generic/cukernel.cu')
-rw-r--r--matrix/generic/cukernel.cu18
1 files changed, 18 insertions, 0 deletions
diff --git a/matrix/generic/cukernel.cu b/matrix/generic/cukernel.cu
index 2e794b7..8b929e4 100644
--- a/matrix/generic/cukernel.cu
+++ b/matrix/generic/cukernel.cu
@@ -113,6 +113,14 @@ __global__ void cudak_(add_row)(const MATRIX_ELEM *a, MATRIX_ELEM *b,
b[j + i * stride] += beta * a[j];
}
+__global__ void cudak_(fill)(MATRIX_ELEM *a,
+ int nrow, int ncol, int stride, double val) {
+ int j = blockIdx.x * blockDim.x + threadIdx.x;
+ int i = blockIdx.y * blockDim.y + threadIdx.y;
+ if (i >= nrow || j >= ncol) return;
+ a[j + i * stride] = val;
+}
+
extern "C" {
#include "../cukernel.h"
@@ -242,5 +250,15 @@ extern "C" {
(MATRIX_ELEM_PTR(a), MATRIX_ELEM_PTR(b), b->nrow, b->ncol,
b->stride / sizeof(MATRIX_ELEM), beta);
}
+
+ void cudak_(cuda_fill)(Matrix *a, double val) {
+ dim3 threadsPerBlock(CUDA_THREADS_N,
+ CUDA_THREADS_N);
+ dim3 numBlocks(CEIL_DIV(a->ncol, threadsPerBlock.x),
+ CEIL_DIV(a->nrow, threadsPerBlock.y));
+ cudak_(fill)<<<numBlocks, threadsPerBlock>>> \
+ (MATRIX_ELEM_PTR(a), a->nrow, a->ncol,
+ a->stride / sizeof(MATRIX_ELEM), val);
+ }
}
#endif