added atomicAdd for select_linear update, however, the result still seems unreproducable, I changed select_linear layer update back to line-by-line

2 files changed, 18 insertions, 2 deletions

diff --git a/nerv/lib/matrix/cukernel.cu b/nerv/lib/matrix/cukernel.cu
index 6fb78f0..1e856b9 100644
--- a/nerv/lib/matrix/cukernel.cu
+++ b/nerv/lib/matrix/cukernel.cu
@@ -2,7 +2,8 @@
 
 #include "cumatrix.h"
 
-__device__ double atomicAdd(double* address, double val) { 
+__device__ double atomicAdd_nvidia(double* address, double val) { 
+    //nvidia provided this implementation on the net
     //atmoicAdd is not included in CUDA for double
     unsigned long long int* address_as_ull =
                                           (unsigned long long int*)address;
@@ -16,6 +17,21 @@ __device__ double atomicAdd(double* address, double val) {
     return __longlong_as_double(old);
 }
 
+__device__ float atomicAdd_nvidia(float* address, float val) { 
+    //nvidia provided this implementation on the net
+    //I tried the included atomocAdd, but the select_liner layer result seems unreproduceable, but sadly, even if I used this implementation, the select_linear layer result is still unreproduceable
+    int* address_as_ull = (int*)address;
+    int old = *address_as_ull, assumed;
+    do {
+        assumed = old;
+        old = atomicCAS(address_as_ull, assumed, 
+                        __float_as_int(val + 
+                        __int_as_float(assumed)));
+    } while (assumed != old);
+    return __int_as_float(old);
+}
+
+
 #define cudak_(NAME) cudak_float_ ## NAME
 #define MATRIX_USE_FLOAT
 #include "generic/elem_type.h"
diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu
index 8885b41..f996fdd 100644
--- a/nerv/lib/matrix/generic/cukernel.cu
+++ b/nerv/lib/matrix/generic/cukernel.cu
@@ -232,7 +232,7 @@ __global__ void cudak_(update_select_rows)(MATRIX_ELEM *c, const MATRIX_ELEM *a,
     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(c + j + i_c * stride_c, c[j + i_c * stride_c] * (- 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,