#include <stdio.h>
#include <math.h>
#include "mytime.h"

#define THREADS 512

// GPU kernel
__global__ void integrate(int *n, double *sum) {
  double h, x;
  int i;

  sum[threadIdx.x] = 0.0;
  h   = 1.0 / (double) *n;
  for (i = threadIdx.x + 1; i <= *n; i += THREADS) {
    x = h * ((double)i - 0.5);
    sum[threadIdx.x] += 4.0 / (1.0 + x*x);
  }
  sum[threadIdx.x] *= h;
}

int main(int argc, char *argv[]) {
  int n, i;
  int *n_d; // device copy of n
  double PI25DT = 3.141592653589793238462643;
  double pi;
  double mypi[THREADS];
  double *mypi_d; // device copy of pi
  struct timeval startwtime, endwtime, diffwtime;
  
  // Allocate memory on GPU
  cudaMalloc( (void **) &n_d, sizeof(int) * 1 );
  cudaMalloc( (void **) &mypi_d, sizeof(double) * THREADS );

  while (1) {
    printf("Enter the number of intervals: (0 quits) ");fflush(stdout);
    scanf("%d",&n);

    gettimeofday(&startwtime, NULL);
    if (n == 0)
      break;

    // copy from CPU to GPU
    cudaMemcpy( n_d, &n, sizeof(int) * 1, cudaMemcpyHostToDevice );

    integrate<<< 1, THREADS>>>(n_d, mypi_d);
    
    // copy back from GPU to CPU
    cudaMemcpy( &mypi, mypi_d, sizeof(double) * THREADS, cudaMemcpyDeviceToHost );

    pi = 0.0;
    for (i = 0; i < THREADS; i++)
      pi += mypi[i];

    gettimeofday(&endwtime, NULL);
    MINUS_UTIME(diffwtime, endwtime, startwtime);
    printf("pi is approximately %.16f, Error is %.16f\n",
	   pi, fabs(pi - PI25DT));
    printf("wall clock time = %d.%06d\n",
	   diffwtime.tv_sec, diffwtime.tv_usec);
  }

  // free GPU memory
  cudaFree(n_d);
  cudaFree(mypi_d);

  return 0;
}