aboutsummaryrefslogtreecommitdiff
path: root/nerv/lib/matrix/generic/cukernel.cu
diff options
context:
space:
mode:
authorDeterminant <ted.sybil@gmail.com>2016-02-15 15:04:13 +0800
committerDeterminant <ted.sybil@gmail.com>2016-02-15 15:04:13 +0800
commit3362020a6bc43766a92882abe6d127c8bb98a628 (patch)
treecad93eb88c2813694c0ae4ca4ecb9873a719ad85 /nerv/lib/matrix/generic/cukernel.cu
parentdcad8a3f80fc55ca93984d981f9b829d2e4ea728 (diff)
try a basic merge
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,