diff options
author | txh18 <cloudygooseg@gmail.com> | 2015-11-17 21:58:05 +0800 |
---|---|---|
committer | txh18 <cloudygooseg@gmail.com> | 2015-11-17 21:58:05 +0800 |
commit | fc4c5a71053b837ed6143659a6e7b45792ed9e51 (patch) | |
tree | a5d3fbf7c59dfb08f67b2156d70def64f98925ad /nerv/lib | |
parent | b8b6bb0a6b9fb9b8d72de42d27f598bfddd1cd0e (diff) |
added atomicAdd for select_linear update, however, the result still seems unreproducable, I changed select_linear layer update back to line-by-line
Diffstat (limited to 'nerv/lib')
-rw-r--r-- | nerv/lib/matrix/cukernel.cu | 18 | ||||
-rw-r--r-- | nerv/lib/matrix/generic/cukernel.cu | 2 |
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, |