aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--nerv/lib/matrix/cukernel.cu1
-rw-r--r--nerv/lib/matrix/cumatrix.c1
-rw-r--r--nerv/lib/matrix/generic/cumatrix.c2
-rw-r--r--nerv/lib/matrix/generic/elem_type.h13
-rw-r--r--nerv/lib/matrix/generic/matrix.c9
-rw-r--r--nerv/lib/matrix/matrix.h1
-rw-r--r--nerv/lib/matrix/mmatrix.c5
-rw-r--r--nerv/matrix/cumatrix.c1
-rw-r--r--nerv/matrix/generic/cumatrix.c2
-rw-r--r--nerv/matrix/generic/elem_type.h22
-rw-r--r--nerv/matrix/generic/mmatrix.c2
-rw-r--r--nerv/matrix/mmatrix.c2
12 files changed, 28 insertions, 33 deletions
diff --git a/nerv/lib/matrix/cukernel.cu b/nerv/lib/matrix/cukernel.cu
index 210e6bf..c20e538 100644
--- a/nerv/lib/matrix/cukernel.cu
+++ b/nerv/lib/matrix/cukernel.cu
@@ -44,6 +44,7 @@ __device__ float atomicAdd_nvidia(float* address, float val) {
#undef MATRIX_USE_FLOAT
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
diff --git a/nerv/lib/matrix/cumatrix.c b/nerv/lib/matrix/cumatrix.c
index ff1168d..a8ed075 100644
--- a/nerv/lib/matrix/cumatrix.c
+++ b/nerv/lib/matrix/cumatrix.c
@@ -57,6 +57,7 @@ void nerv_cumatrix_init() {
#undef MATRIX_USE_FLOAT
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
#undef MATRIX_CUMATRIX_HOST_TNAME
diff --git a/nerv/lib/matrix/generic/cumatrix.c b/nerv/lib/matrix/generic/cumatrix.c
index 00af895..7643c01 100644
--- a/nerv/lib/matrix/generic/cumatrix.c
+++ b/nerv/lib/matrix/generic/cumatrix.c
@@ -315,7 +315,7 @@ void nerv_matrix_(copy_rows_fromh_by_idx)(Matrix *a, const Matrix *b,
long nrow = a->nrow;
if (!(0 <= b_begin && b_begin + nrow <= idx->ncol))
NERV_EXIT_STATUS(status, MAT_INVALID_COPY_INTERVAL, 0);
- float *idx_ptr = idx->data.f;
+ float *idx_ptr = MATRIX_ELEM_PTR_F(idx);
int i;
if (idx->nrow != 1)
NERV_EXIT_STATUS(status, MAT_IDX_VECTOR_EXP, 0);
diff --git a/nerv/lib/matrix/generic/elem_type.h b/nerv/lib/matrix/generic/elem_type.h
index bffe940..07f6355 100644
--- a/nerv/lib/matrix/generic/elem_type.h
+++ b/nerv/lib/matrix/generic/elem_type.h
@@ -1,22 +1,29 @@
+#define MATRIX_ELEM_PTR_F(self) ((float *)((char *)((self)->data.f) + (self)->offset))
+#define MATRIX_ELEM_PTR_D(self) ((double *)((char *)((self)->data.d) + (self)->offset))
+#define MATRIX_ELEM_PTR_I(self) ((long *)((char *)((self)->data.i) + (self)->offset))
+
#ifdef MATRIX_USE_FLOAT
#define MATRIX_ELEM float
#define MATRIX_ELEM_FMT "%f"
#define MATRIX_ELEM_WRITE_FMT "%.8f"
-#define MATRIX_ELEM_PTR(self) ((self)->data.f)
+#define MATRIX_ELEM_PTR(self) MATRIX_ELEM_PTR_F(self)
+#define MATRIX_ELEM_PTR_BASE(self) ((self)->data.f)
#elif defined(MATRIX_USE_DOUBLE)
#define MATRIX_ELEM double
#define MATRIX_ELEM_FMT "%lf"
#define MATRIX_ELEM_WRITE_FMT "%.8lf"
-#define MATRIX_ELEM_PTR(self) ((self)->data.d)
+#define MATRIX_ELEM_PTR(self) MATRIX_ELEM_PTR_D(self)
+#define MATRIX_ELEM_PTR_BASE(self) ((self)->data.d)
#elif defined(MATRIX_USE_INT)
#define MATRIX_ELEM long
#define MATRIX_ELEM_FMT "%ld"
#define MATRIX_ELEM_WRITE_FMT "%ld"
-#define MATRIX_ELEM_PTR(self) ((self)->data.i)
+#define MATRIX_ELEM_PTR(self) MATRIX_ELEM_PTR_I(self)
+#define MATRIX_ELEM_PTR_BASE(self) ((self)->data.i)
#endif
diff --git a/nerv/lib/matrix/generic/matrix.c b/nerv/lib/matrix/generic/matrix.c
index 4246751..998d107 100644
--- a/nerv/lib/matrix/generic/matrix.c
+++ b/nerv/lib/matrix/generic/matrix.c
@@ -4,12 +4,11 @@
/* FIXME: malloc failure detection */
void nerv_matrix_(data_free)(Matrix *self, Status *status) {
- if(*self->data_ref == 0) return; /* FIXME: repeat free memory */
assert(*self->data_ref > 0);
if (--(*self->data_ref) == 0)
{
/* free matrix data */
- MATRIX_DATA_FREE(MATRIX_ELEM_PTR(self), status);
+ MATRIX_DATA_FREE(MATRIX_ELEM_PTR_BASE(self), status);
free(self->data_ref);
free(self);
}
@@ -29,7 +28,7 @@ Matrix *nerv_matrix_(create)(long nrow, long ncol, Status *status) {
self->ncol = ncol;
self->nmax = self->nrow * self->ncol;
self->dim = 2;
- MATRIX_DATA_ALLOC(&MATRIX_ELEM_PTR(self), &self->stride,
+ MATRIX_DATA_ALLOC(&MATRIX_ELEM_PTR_BASE(self), &self->stride,
sizeof(MATRIX_ELEM) * self->ncol, self->nrow,
status);
if (status->err_code != NERV_NORMAL)
@@ -39,6 +38,7 @@ Matrix *nerv_matrix_(create)(long nrow, long ncol, Status *status) {
}
self->data_ref = (long *)malloc(sizeof(long));
*self->data_ref = 0;
+ self->offset = 0;
nerv_matrix_(data_retain)(self);
NERV_SET_STATUS(status, NERV_NORMAL, 0);
return self;
@@ -55,8 +55,9 @@ Matrix *nerv_matrix_(getrow)(Matrix *self, int row) {
prow->dim = 1;
prow->stride = self->stride;
prow->nmax = prow->ncol;
- MATRIX_ELEM_PTR(prow) = MATRIX_ROW_PTR(self, row);
+ prow->data = self->data;
prow->data_ref = self->data_ref;
+ prow->offset = row * self->stride;
nerv_matrix_(data_retain)(prow);
return prow;
}
diff --git a/nerv/lib/matrix/matrix.h b/nerv/lib/matrix/matrix.h
index 67a6e30..073bd13 100644
--- a/nerv/lib/matrix/matrix.h
+++ b/nerv/lib/matrix/matrix.h
@@ -12,6 +12,7 @@ typedef struct Matrix {
double *d;
long *i;
} data; /* pointer to actual storage */
+ unsigned long offset; /* the actual beginning of the matrix */
long *data_ref;
} Matrix;
diff --git a/nerv/lib/matrix/mmatrix.c b/nerv/lib/matrix/mmatrix.c
index b8157eb..b5670f2 100644
--- a/nerv/lib/matrix/mmatrix.c
+++ b/nerv/lib/matrix/mmatrix.c
@@ -6,6 +6,7 @@
#define host_matrix_(NAME) host_matrix_float_##NAME
#define nerv_matrix_(NAME) nerv_matrix_host_float_##NAME
#include "generic/matrix.h"
+#include "generic/elem_type.h"
#include "generic/mmatrix.c"
Matrix *nerv_matrix_(perm_gen)(int ncol, Status *status) {
@@ -13,7 +14,7 @@ Matrix *nerv_matrix_(perm_gen)(int ncol, Status *status) {
Matrix *self = nerv_matrix_(create)(1, ncol, status);
if (status->err_code != NERV_NORMAL)
return NULL;
- float *prow = self->data.f;
+ float *prow = MATRIX_ELEM_PTR_F(self);
for (i = 0; i < ncol; i++)
prow[i] = i;
for (i = ncol - 1; i >= 0; i--)
@@ -31,6 +32,7 @@ Matrix *nerv_matrix_(perm_gen)(int ncol, Status *status) {
#undef MATRIX_USE_FLOAT
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
@@ -44,6 +46,7 @@ Matrix *nerv_matrix_(perm_gen)(int ncol, Status *status) {
#undef MATRIX_USE_DOUBLE
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
diff --git a/nerv/matrix/cumatrix.c b/nerv/matrix/cumatrix.c
index fef03fc..bf92f92 100644
--- a/nerv/matrix/cumatrix.c
+++ b/nerv/matrix/cumatrix.c
@@ -49,6 +49,7 @@ const char *nerv_matrix_(tname) = "nerv.CuMatrixFloat";
#undef MATRIX_USE_FLOAT
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
#undef MATRIX_CUMATRIX_HOST_TNAME
diff --git a/nerv/matrix/generic/cumatrix.c b/nerv/matrix/generic/cumatrix.c
index e1519b0..be3d627 100644
--- a/nerv/matrix/generic/cumatrix.c
+++ b/nerv/matrix/generic/cumatrix.c
@@ -1,5 +1,5 @@
#ifdef NERV_GENERIC_CUMATRIX
-#include "elem_type.h"
+#include "../../lib/matrix/generic/elem_type.h"
#define MATRIX_DATA_WRITE(L, data, idx, val) cuda_matrix_(write)(L, data, idx, val)
#define MATRIX_DATA_READ(L, data, idx) cuda_matrix_(read)(L, data, idx)
#define MATRIX_INIT(L) cuda_matrix_(init)(L)
diff --git a/nerv/matrix/generic/elem_type.h b/nerv/matrix/generic/elem_type.h
deleted file mode 100644
index bffe940..0000000
--- a/nerv/matrix/generic/elem_type.h
+++ /dev/null
@@ -1,22 +0,0 @@
-#ifdef MATRIX_USE_FLOAT
-
-#define MATRIX_ELEM float
-#define MATRIX_ELEM_FMT "%f"
-#define MATRIX_ELEM_WRITE_FMT "%.8f"
-#define MATRIX_ELEM_PTR(self) ((self)->data.f)
-
-#elif defined(MATRIX_USE_DOUBLE)
-
-#define MATRIX_ELEM double
-#define MATRIX_ELEM_FMT "%lf"
-#define MATRIX_ELEM_WRITE_FMT "%.8lf"
-#define MATRIX_ELEM_PTR(self) ((self)->data.d)
-
-#elif defined(MATRIX_USE_INT)
-
-#define MATRIX_ELEM long
-#define MATRIX_ELEM_FMT "%ld"
-#define MATRIX_ELEM_WRITE_FMT "%ld"
-#define MATRIX_ELEM_PTR(self) ((self)->data.i)
-
-#endif
diff --git a/nerv/matrix/generic/mmatrix.c b/nerv/matrix/generic/mmatrix.c
index 01dd9e5..a4e8489 100644
--- a/nerv/matrix/generic/mmatrix.c
+++ b/nerv/matrix/generic/mmatrix.c
@@ -1,6 +1,6 @@
#ifdef NERV_GENERIC_MMATRIX
#include "../../lib/matrix/generic/matrix.h"
-#include "elem_type.h"
+#include "../../lib/matrix/generic/elem_type.h"
#define MATRIX_DATA_WRITE(L, data, idx, val) (data[idx] = val)
#define MATRIX_DATA_READ(L, data, idx) (data[idx])
#define MATRIX_INIT(L) host_matrix_(init)(L)
diff --git a/nerv/matrix/mmatrix.c b/nerv/matrix/mmatrix.c
index 961059c..20c31d6 100644
--- a/nerv/matrix/mmatrix.c
+++ b/nerv/matrix/mmatrix.c
@@ -45,6 +45,7 @@ static const luaL_Reg nerv_matrix_(extra_methods_int)[] = {
#undef MATRIX_USE_FLOAT
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
#undef MMATRIX_INIT
@@ -60,6 +61,7 @@ const char *nerv_matrix_(tname) = "nerv.MMatrixDouble";
#undef MATRIX_USE_DOUBLE
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
+#undef MATRIX_ELEM_PTR_BASE
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT
Powered by cgit, Themed by Ted Yin.