         //    }
         //}
     }
-	//Flush FPEs to superaccs
+    //Flush FPEs to superaccs
     #pragma unroll
     for(uint i = 0; i != NBFPE; ++i)
         Accumulate(l_workingBase, a[i]);

            sum += d_PartialSuperaccs[i * BIN_COUNT + lid];

        d_PartialSuperaccs[lid] = sum;

        //__syncthreads();
        //if (lid == 0)
        //    d_Res[0] = Round(d_PartialSuperaccs);
    }
}

__host__
std::vector<int64_t> exdot_gpu(unsigned size, const double* x1_ptr, const double* x2_ptr)
{
    thrust::device_vector<long long int> d_PartialSuperaccsV( PARTIAL_SUPERACCS_COUNT*BIN_COUNT); //39 columns and PSC rows
    long long int *d_PartialSuperaccs = thrust::raw_pointer_cast( d_PartialSuperaccsV.data());
    ExDOT<<<PARTIAL_SUPERACCS_COUNT, WORKGROUP_SIZE>>>( d_PartialSuperaccs, x1_ptr, x2_ptr,size);
    thrust::device_vector<double> r(1,0);
    double *r_ptr = thrust::raw_pointer_cast( r.data());
    ExDOTComplete<<<PARTIAL_SUPERACCS_COUNT/MERGE_SUPERACCS_SIZE, MERGE_WORKGROUP_SIZE>>>( r_ptr, d_PartialSuperaccs );
    std::vector<int64_t> h_Superacc(BIN_COUNT);
    cudaMemcpy( &h_Superacc[0], d_PartialSuperaccs, BIN_COUNT*sizeof(long long int), cudaMemcpyDeviceToHost);
    return h_Superacc;

////////////////////////////////////////////////////////////////////////////////
__global__
void ExDOTComplete(
     long long int *d_PartialSuperaccs
) {
    uint lid = threadIdx.x;
