diff options
| -rw-r--r-- | Makefile | 6 | ||||
| -rw-r--r-- | nerv/lib/matrix/generic/cukernel.cu | 8 |
2 files changed, 11 insertions, 3 deletions
@@ -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]; } |