aboutsummaryrefslogtreecommitdiff
path: root/nerv/lib/matrix/cukernel.cu
diff options
context:
space:
mode:
Diffstat (limited to 'nerv/lib/matrix/cukernel.cu')
-rw-r--r--nerv/lib/matrix/cukernel.cu18
1 files changed, 17 insertions, 1 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"