     for( unsigned i=0; i<10; i++)
         std::cout << d_PartialSuperaccsV[i]<<" ";
     cudaDeviceSynchronize();
-    return result[0];
+    double sum = result[0];
+    return sum;
 }
 }//namespace exblas
 

    *s = (a - (r - z)) + (b - z);
    return r;
}
// signedcarry in {-1, 0, 1}
__device__ 
INTEGER xadd( INTEGER *sa, INTEGER x, unsigned char *of) {
    // OF and SF  -> carry=1
    // OF and !SF -> carry=-1
    // !OF        -> carry=0
    //INTEGER y = atom_add(sa, x);
    INTEGER y = atomicAdd((unsigned long long*)sa, (unsigned long long)x); //not sure if this is the correct CUDA function
    INTEGER z = y + x; // since the value sa->superacc[i] can be changed by another work item

    // TODO: cover also underflow


namespace exblas{

typedef long long int INTEGER;

static constexpr uint BIN_COUNT      =  39;
static constexpr uint KRX            =  8;                 // High-radix carry-save bits
