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


namespace exblas{

typedef long INTEGER;

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

    *s = (a - (r - z)) + (b - z);
    return r;
}
__device__ INTEGER atomicAdd( INTEGER* address, INTEGER val)
{
    unsigned long long int* address_as_ull =
        (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull;

    old = atomicCAS(address_as_ull, old,
                      (unsigned long long int)(val + (INTEGER)old));
    //assume that bit patterns don't change when casting
    return (INTEGER)(old);
}
// 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(sa, x); 
    INTEGER z = y + x; // since the value sa->superacc[i] can be changed by another work item

    // TODO: cover also underflow
