aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--Makefile5
-rw-r--r--cumatrix_example.lua1
-rw-r--r--matrix/cukernel.cu21
-rw-r--r--matrix/cukernel.h4
-rw-r--r--matrix/cumatrix.c10
5 files changed, 40 insertions, 1 deletions
diff --git a/Makefile b/Makefile
index 4514263..bfda339 100644
--- a/Makefile
+++ b/Makefile
@@ -1,5 +1,5 @@
.PHONY: all clean luajit
-OBJS := oop_example.o nerv.o luaT.o common.o matrix/matrix.o matrix/cumatrix.o matrix/init.o
+OBJS := oop_example.o nerv.o luaT.o common.o matrix/matrix.o matrix/cumatrix.o matrix/init.o matrix/cukernel.o
LIBS := libnerv.so
LUA_LIBS := matrix/init.lua nerv.lua
INCLUDE := -I build/luajit-2.0/include/luajit-2.0/ -DLUA_USE_APICHECK
@@ -10,6 +10,7 @@ LDFLAGS := -L$(CUDA_BASE)/lib64/ -Wl,-rpath=$(CUDA_BASE)/lib64/ -lcudart -lcubl
CFLAGS :=
OBJ_DIR := build/objs
LUA_DIR := build/lua
+NVCC := $(CUDA_BASE)/bin/nvcc
OBJS := $(addprefix $(OBJ_DIR)/,$(OBJS))
LIBS := $(addprefix $(OBJ_DIR)/,$(LIBS))
@@ -28,6 +29,8 @@ $(OBJ_DIR)/%.o: %.c
gcc -c -o $@ $< $(INCLUDE) -fPIC $(CFLAGS)
$(OBJ_DIR)/matrix/%.o: matrix/%.c
gcc -c -o $@ $< $(INCLUDE) -fPIC $(CFLAGS)
+$(OBJ_DIR)/matrix/cukernel.o: matrix/cukernel.cu
+ $(NVCC) -c -o $@ $< -Xcompiler -fPIC $(INCLUDE) $(CFLAGS)
$(LUA_DIR)/%.lua: %.lua
cp $< $@
$(OBJ_DIR)/luaT.o:
diff --git a/cumatrix_example.lua b/cumatrix_example.lua
index ccd88b8..ce11eea 100644
--- a/cumatrix_example.lua
+++ b/cumatrix_example.lua
@@ -25,3 +25,4 @@ d[1][1] = 2
d[2][2] = 3
print(d)
print(t * d)
+print(t:sigmoid())
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));
+}
diff --git a/matrix/cukernel.h b/matrix/cukernel.h
new file mode 100644
index 0000000..5b9e3a6
--- /dev/null
+++ b/matrix/cukernel.h
@@ -0,0 +1,4 @@
+#ifndef NERV_CUKERNEL_H
+#define NERV_CUKERNEL_H
+void cuda_sigmoid(const Matrix *a, Matrix *b);
+#endif
diff --git a/matrix/cumatrix.c b/matrix/cumatrix.c
index 9c2878a..7759ca1 100644
--- a/matrix/cumatrix.c
+++ b/matrix/cumatrix.c
@@ -7,6 +7,7 @@
#define nerv_float_matrix_(NAME) nerv_float_matrix_cuda_ ## NAME
#include "../common.h"
#include "generic/matrix.h"
+#include "cukernel.h"
#include "cuda.h"
#include "driver_types.h"
#include "cublas_v2.h"
@@ -56,9 +57,18 @@ static int nerv_float_matrix_(mul)(lua_State *L) {
return 1;
}
+static int nerv_float_matrix_(sigmoid)(lua_State *L) {
+ Matrix *a = luaT_checkudata(L, 1, nerv_float_matrix_(tname));
+ Matrix *b = nerv_float_matrix_(new_)(a->nrow, a->ncol);
+ cuda_sigmoid(a, b);
+ luaT_pushudata(L, b, nerv_float_matrix_(tname));
+ return 1;
+}
+
static const luaL_Reg nerv_float_matrix_(extra_methods)[] = {
{"__add__", nerv_float_matrix_(add)},
{"__mul__", nerv_float_matrix_(mul)},
+ {"sigmoid", nerv_float_matrix_(sigmoid)},
{NULL, NULL}
};