From c640c9d33df54426c0ff66d242475cfb11ab86c5 Mon Sep 17 00:00:00 2001 From: Determinant Date: Sat, 12 Mar 2016 14:12:50 +0800 Subject: ready for alpha-2 release --- nerv/lib/matrix/generic/cukernel.cu | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) (limited to 'nerv') 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]; } -- cgit v1.2.3