aboutsummaryrefslogtreecommitdiff
path: root/nerv/lib/matrix/generic/cukernel.cu
diff options
context:
space:
mode:
Diffstat (limited to 'nerv/lib/matrix/generic/cukernel.cu')
-rw-r--r--nerv/lib/matrix/generic/cukernel.cu6
1 files changed, 6 insertions, 0 deletions
diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu
index 2b696d5..7f780a8 100644
--- a/nerv/lib/matrix/generic/cukernel.cu
+++ b/nerv/lib/matrix/generic/cukernel.cu
@@ -356,11 +356,17 @@ __global__ void cudak_(copy_rows_by_idx)(const MATRIX_ELEM *a, MATRIX_ELEM *b,
int j = blockIdx.x * blockDim.x + threadIdx.x;
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= nrow || j >= ncol) return;
+ /*
int k = lrintf(idx[i]);
if (k < 0 || k >= a_nrow) {
printf("error in kernel copy_rows_by_idx k(%d) out of range\n", k);
}
b[j + i * stride] = a[j + k * stride];
+ */
+ /* NOTE: in most cases it is guaranteed
+ * the idx is within the range, checking
+ * would bring some overhead. */
+ b[j + i * stride] = a[j + lrintf(idx[i]) * stride];
}
__global__ void cudak_(copy_rows_by_colidx)(const MATRIX_ELEM *a, MATRIX_ELEM *b,