     double mine = xs[idx] * ys[idx];
     atomicAdd(&g_res, mine);
 }
+
+inline void gpuAssert(cudaError_t err, const char* file, int line) {
+    if( err != cudaSuccess ) {
+        fprintf(stderr,"GPU error: %s %s %d\n", cudaGetErrorString(err), file, line);
+        exit(1);
+    }
+}
+#define GPU_ASSERT(ans) {gpuAssert(ans,__FILE__,__LINE__);}
 
 /* Execute kernel on GPU */
 extern "C"

    double *buf_xs = NULL;
    double *buf_ys = NULL;
    // Allocate memory and copy to device
    cudaMalloc( &buf_xs, n*sizeof(double) );
    cudaMalloc( &buf_ys, n*sizeof(double) );
    cudaMemcpy(buf_xs, xs, n*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(buf_ys, ys, n*sizeof(double), cudaMemcpyHostToDevice);
    // Execute kernel
    dotp_kernel<<<grid, block>>>(buf_xs, buf_ys);
    // Free buffers & obtain data
    cudaFree( &buf_xs );
    cudaFree( &buf_ys );
    double res;
    cudaMemcpyFromSymbol(&res, g_res, sizeof(double), 0, cudaMemcpyDeviceToHost);
    return res;
}


#include <stdio.h>
/* Global variable for returning result of computation */
__device__ float g_res;

