aboutsummaryrefslogtreecommitdiff
path: root/nerv
diff options
context:
space:
mode:
authortxh18 <[email protected]>2015-11-18 14:59:42 +0800
committertxh18 <[email protected]>2015-11-18 14:59:42 +0800
commit22dac37c663605aa6c6fa0426696d2d01da4370f (patch)
treeba4c3d66d69361e4f6343e5be8f69a09ae94a07a /nerv
parent0acdf7bd21a1953d8e5ba07fb39a7187f049527f (diff)
small coding style change
Diffstat (limited to 'nerv')
-rw-r--r--nerv/lib/matrix/generic/cukernel.cu5
1 files changed, 3 insertions, 2 deletions
diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu
index f996fdd..e1063af 100644
--- a/nerv/lib/matrix/generic/cukernel.cu
+++ b/nerv/lib/matrix/generic/cukernel.cu
@@ -231,8 +231,9 @@ __global__ void cudak_(update_select_rows)(MATRIX_ELEM *c, const MATRIX_ELEM *a,
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= nrow_a || j >= ncol_a) return;
int i_c = lrintf(idx[i]);
- //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);
+ //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);
}
__global__ void cudak_(expand_frm)(const MATRIX_ELEM *a, MATRIX_ELEM *b,