From 0f953414dfdbd7abb7b867ce0c3f9390551c1083 Mon Sep 17 00:00:00 2001
From: Determinant <ted.sybil@gmail.com>
Date: Mon, 18 May 2015 11:09:58 +0800
Subject: add sigmoid for cumatrix

---
 matrix/cukernel.cu | 21 +++++++++++++++++++++
 1 file changed, 21 insertions(+)
 create mode 100644 matrix/cukernel.cu

(limited to 'matrix/cukernel.cu')

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));
+}
-- 
cgit v1.2.3