aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorDeterminant <ted.sybil@gmail.com>2016-03-12 14:12:50 +0800
committerDeterminant <ted.sybil@gmail.com>2016-03-12 14:12:50 +0800
commitc640c9d33df54426c0ff66d242475cfb11ab86c5 (patch)
tree17f6f92100f480ef2c0af6a133589e6cbb8b865a
parent2b03555ea53a47e87d03a79feb866c868d424f01 (diff)
ready for alpha-2 release
-rw-r--r--Makefile6
-rw-r--r--nerv/lib/matrix/generic/cukernel.cu8
2 files changed, 11 insertions, 3 deletions
diff --git a/Makefile b/Makefile
index 28012da..728d14b 100644
--- a/Makefile
+++ b/Makefile
@@ -23,7 +23,7 @@ export KALDI_BASE
export BLAS_LDFLAGS
.PHONY: nerv speech/speech_utils speech/htk_io speech/kaldi_io speech/kaldi_decode \
- nerv-clean speech/speech_utils-clean speech/htk_io-clean speech/kaldi_io-clean speech/kaldi_decode-clean \
+ nerv-clean speech-clean speech/speech_utils-clean speech/htk_io-clean speech/kaldi_io-clean speech/kaldi_decode-clean \
Penlight
all: luajit luarocks Penlight nerv
@@ -35,8 +35,10 @@ luarocks:
speech: speech/speech_utils speech/htk_io speech/kaldi_io speech/kaldi_decode
speech-clean: speech/speech_utils-clean speech/htk_io-clean speech/kaldi_io-clean speech/kaldi_decode-clean
clean: nerv-clean speech-clean
+uninstall:
+ -rm -rf install/
nerv Penlight speech/speech_utils speech/htk_io speech/kaldi_io speech/kaldi_decode:
cd $@; $(PREFIX)/bin/luarocks make
nerv-clean speech/speech_utils-clean speech/htk_io-clean speech/kaldi_io-clean speech/kaldi_decode-clean:
- cd $(subst -clean,,$@); make clean LUA_BINDIR=$(PREFIX)/bin/
+ -make -C $(subst -clean,,$@) clean LUA_BINDIR=$(PREFIX)/bin/
diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu
index 93121dc..fc630ad 100644
--- a/nerv/lib/matrix/generic/cukernel.cu
+++ b/nerv/lib/matrix/generic/cukernel.cu
@@ -277,11 +277,13 @@ __global__ void cudak_(update_select_rows_by_rowidx)(MATRIX_ELEM *c, const MATRI
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= nrow_a || j >= ncol_a) return;
int i_c = lrintf(idx[i]);
+ /*
if (i_c < 0 || i_c >= nrow_c) {
printf("ERROR inside kernel update_select_rows, i_c(%d) out of range!", i_c);
}
+ */
//critical: i_c could conflict among threads(same index in the idx array), so atomicAdd is used
- //c[j + i_c * stride_c] = c[j + i_c * stride_c] * (1 - beta * alpha) + a[j + i * stride_a] * alpha;
+ //c[j + i_c * stride_c] = c[j + i_c * stride_c] * (1 - beta * alpha) + a[j + i * stride_a] * alpha;
atomicAdd_nvidia(c + j + i_c * stride_c, c[j + i_c * stride_c] * (- beta * alpha) + a[j + i * stride_a] * alpha);
}
@@ -291,9 +293,11 @@ __global__ void cudak_(update_select_rows_by_colidx)(MATRIX_ELEM *c, const MATRI
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= nrow_a || j >= ncol_a) return;
int i_c = lrintf(idx[stride_idx * i]);
+ /*
if (i_c < 0 || i_c >= nrow_c) {
printf("ERROR inside kernel update_select_rows, i_c(%d) out of range!", i_c);
}
+ */
//critical: i_c could conflict among threads(same index in the idx array), so atomicAdd is used
//c[j + i_c * stride_c] = c[j + i_c * stride_c] * (1 - beta * alpha) + a[j + i * stride_a] * alpha;
atomicAdd_nvidia(c + j + i_c * stride_c, c[j + i_c * stride_c] * (- beta * alpha) + a[j + i * stride_a] * alpha);
@@ -385,9 +389,11 @@ __global__ void cudak_(copy_rows_by_colidx)(const MATRIX_ELEM *a, MATRIX_ELEM *b
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= nrow || j >= ncol) return;
int k = lrintf(idx[i * idx_stride]);
+ /*
if (k < 0 || k >= a_nrow) {
printf("error in kernel copy_rows_by_colidx k(%d) out of range\n", k);
}
+ */
b[j + i * stride] = a[j + k * stride];
}