From 3362020a6bc43766a92882abe6d127c8bb98a628 Mon Sep 17 00:00:00 2001 From: Determinant Date: Mon, 15 Feb 2016 15:04:13 +0800 Subject: try a basic merge --- nerv/lib/matrix/generic/cukernel.cu | 6 ++++++ 1 file changed, 6 insertions(+) (limited to 'nerv/lib/matrix/generic/cukernel.cu') 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, -- cgit v1.2.3