diff options
| -rw-r--r-- | Makefile | 5 | ||||
| -rw-r--r-- | cumatrix_example.lua | 1 | ||||
| -rw-r--r-- | matrix/cukernel.cu | 21 | ||||
| -rw-r--r-- | matrix/cukernel.h | 4 | ||||
| -rw-r--r-- | matrix/cumatrix.c | 10 |
5 files changed, 40 insertions, 1 deletions
@@ -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} }; |