aboutsummaryrefslogtreecommitdiff
path: root/matrix/cukernel.cu
diff options
context:
space:
mode:
authorDeterminant <[email protected]>2015-05-18 11:09:58 +0800
committerDeterminant <[email protected]>2015-05-18 11:09:58 +0800
commit0f953414dfdbd7abb7b867ce0c3f9390551c1083 (patch)
tree7b50dd77c0316acac8cdd8c83ad5a34b99b31a83 /matrix/cukernel.cu
parent83006367aeec856bf8e59231c78df5b1802e3138 (diff)
add sigmoid for cumatrix
Diffstat (limited to 'matrix/cukernel.cu')
-rw-r--r--matrix/cukernel.cu21
1 files changed, 21 insertions, 0 deletions
diff --git a/matrix/cukernel.cu b/matrix/cukernel.cu
new file mode 100644
index 0000000..91e7e35
--- /dev/null
+++ b/matrix/cukernel.cu
@@ -0,0 +1,21 @@
+#include "generic/matrix.h"
+#define CUDA_THREADS_N 16
+#define CEIL_DIV(a, b) (((a) + (b) - 1) / (b))
+__global__ void sigmoid(const float *a, float *b,
+ int nrow, int ncol, int stride) {
+ int j = blockIdx.x * blockDim.x + threadIdx.x;
+ int i = blockIdx.y * blockDim.y + threadIdx.y;
+ long idx;
+ if (i >= nrow || j >= ncol) return;
+ idx = j + i * stride;
+ b[idx] = 1.0 / (1.0 + exp(-a[idx]));
+}
+
+extern "C" void cuda_sigmoid(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));
+ sigmoid<<<numBlocks, threadsPerBlock>>>(a->data.f, b->data.f, b->nrow, b->ncol,
+ b->stride / sizeof(float));
+}