path: root/nerv/lib/matrix/cukernel.cu

#define NERV_GENERIC_CUKERNEL

#include "cumatrix.h"

__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;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, 
                        __double_as_longlong(val + 
                        __longlong_as_double(assumed)));
    } while (assumed != old);
    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"
#include "generic/cukernel.cu"
#undef cudak_
#undef MATRIX_USE_FLOAT
#undef MATRIX_ELEM
#undef MATRIX_ELEM_PTR
#undef MATRIX_ELEM_FMT
#undef MATRIX_ELEM_WRITE_FMT

#define cudak_(NAME) cudak_double_ ## NAME
#define MATRIX_USE_DOUBLE
#include "generic/elem_type.h"
#include "generic/cukernel.cu"