#include <deque>
#include <vector>
#include <fstream>
#include <iostream>

using namespace std;

#include "ini.cxx"
#include "pro.cu"

dats e=d; // copy of "d" for device
dats *ed; // pointer to structure on device

unsigned int pmax, pn;
photon * pz;

float deviceTime=0;
cudaDeviceProp prop;
int nblk;

void checkError(cudaError result){
  if(result!=cudaSuccess){
    cerr<<"Error: "<<cudaGetErrorString(result)<<endl;
    exit(1);
  }
}

void ini(int type){
  {
    nblk=type==0?prop.multiProcessorCount:WNUM;
    pmax=nblk*NTHR*NPHO;
  }

  {
    pn=0; e.type=type; e.hidx=0;
    for(int m=0; m<3; m++) e.r[m]=0;
  }

  {
    unsigned int size=d.rsize;
    if(size<nblk*NTHR) cerr<<"Error: not enough multipliers: only have "<<size<<" (need "<<(nblk*NTHR)<<")!"<<endl;
  }

  {
    unsigned int size=d.rsize*sizeof(unsigned int);
    checkError(cudaMalloc((void**) &e.rm, size));
    checkError(cudaMemcpy(e.rm, d.rm, size, cudaMemcpyHostToDevice));
  }

  {
    unsigned int size=d.rsize*sizeof(unsigned long long);
    checkError(cudaMalloc((void**) &e.rs, size));
    checkError(cudaMemcpy(e.rs, d.rs, size, cudaMemcpyHostToDevice));
  }

  for(int i=0; i<WNUM; i++){
    unsigned int size=sizeof(ices);
    checkError(cudaMalloc((void**) &e.w[i], size));
    checkError(cudaMemcpy(e.w[i], d.w[i], size, cudaMemcpyHostToDevice));
  }

  {
    unsigned int size=d.hnum*sizeof(hit);
    checkError(cudaMalloc((void**) &e.hits, size));
  }

  {
    pz = new photon[pmax];
    unsigned int size=pmax*sizeof(photon);
    checkError(cudaMalloc((void**) &e.pz, size));
  }

  {
    unsigned int size=d.gsize;
    checkError(cudaMemcpyToSymbol(oms, d.oms, size*sizeof(DOM)));
  }

  {
    checkError(cudaMalloc((void**) &ed, sizeof(dats)));
    checkError(cudaMemcpy(ed, &e, sizeof(dats), cudaMemcpyHostToDevice));
  }
}

void fin(){
  checkError(cudaFree(e.rm));
  checkError(cudaFree(e.rs));
  for(int i=0; i<WNUM; i++) checkError(cudaFree(e.w[i]));
  checkError(cudaFree(e.hits));
  checkError(cudaFree(e.pz));
  checkError(cudaFree(ed));
  delete pz;
}

void listDevices(){
  int deviceCount, driver, runtime;
  cudaGetDeviceCount(&deviceCount);
  cudaDriverGetVersion(&driver);
  cudaRuntimeGetVersion(&runtime);
  fprintf(stderr, "\nFound %d devices, driver %d, runtime %d\n", deviceCount, driver, runtime);
  for(int device=0; device<deviceCount; ++device){
    cudaGetDeviceProperties(&prop, device);
    fprintf(stderr, "%d(%d.%d): %s %g GHz G(%lu) S(%lu) C(%lu) R(%d) W(%d)\n"
	    "\tl%d o%d c%d h%d i%d m%d a%lu M(%lu) T(%d: %d,%d,%d) G(%d,%d,%d)\n",
	    device, prop.major, prop.minor, prop.name, prop.clockRate/1.e6,
	    (unsigned long)prop.totalGlobalMem, (unsigned long)prop.sharedMemPerBlock,
	    (unsigned long)prop.totalConstMem, prop.regsPerBlock, prop.warpSize,
	    prop.kernelExecTimeoutEnabled, prop.deviceOverlap, prop.computeMode,
	    prop.canMapHostMemory, prop.integrated, prop.multiProcessorCount,
	    (unsigned long)prop.textureAlignment,
	    (unsigned long)prop.memPitch, prop.maxThreadsPerBlock,
	    prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2],
	    prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  }
}

#include "f2k.cxx"

void kernel(unsigned int num){
  {
    float dt;
    cudaEvent_t evt1, evt2;

    checkError(cudaEventCreateWithFlags(&evt1, cudaEventBlockingSync));
    checkError(cudaEventCreateWithFlags(&evt2, cudaEventBlockingSync));

    checkError(cudaEventRecord(evt1, NULL));
    checkError(cudaEventSynchronize(evt1));

    propagate<<< nblk, NTHR >>>(ed, num);
    checkError(cudaGetLastError());

    checkError(cudaEventRecord(evt2, NULL));
    checkError(cudaEventSynchronize(evt2));

    checkError(cudaEventElapsedTime(&dt, evt1, evt2));
    deviceTime+=dt;

    checkError(cudaEventDestroy(evt1));
    checkError(cudaEventDestroy(evt2));

    checkError(cudaThreadSynchronize());
  }

  {
    checkError(cudaMemcpy(&e, ed, sizeof(int), cudaMemcpyDeviceToHost));
    unsigned int size=e.hidx*sizeof(hit);
    checkError(cudaMemcpy(d.hits, e.hits, size, cudaMemcpyDeviceToHost));
  }

  cerr<<"photons: "<<num<<"  hits: "<<e.hidx<<endl;

  print();

  if(e.hidx>=e.hnum) cerr<<"Error: data buffer overflow occurred!"<<endl;
}

void flasher(int str, int dom, unsigned long long num, int itr){
  ini(0);

  e.hidx=0;
  for(int n=0; n<d.gsize; n++) if(d.names[n].str==str && d.names[n].dom==dom){ for(int m=0; m<3; m++) e.r[m]=d.oms[n].r[m]; break; }
  checkError(cudaMemcpy(ed, &e, 4*sizeof(int), cudaMemcpyHostToDevice));

  for(int j=0; j<max(1, itr); j++){
    for(long long i=num; i>0; i-=pmax){
      if(i<num){
	e.hidx=0;
	checkError(cudaMemcpy(ed, &e, sizeof(int), cudaMemcpyHostToDevice));
      }

      kernel(min(i, (unsigned long long) pmax));
    }
    if(itr>0) printf("\n");
  }

  fin();
}

void output(){
  {
    e.hidx=0;
    checkError(cudaMemcpy(ed, &e, sizeof(int), cudaMemcpyHostToDevice));
  }

  {
    unsigned int size=pn*sizeof(photon);
    checkError(cudaMemcpy(e.pz, pz, size, cudaMemcpyHostToDevice));
  }

  kernel(pn);

  pn=0;
}

int main(int arg_c, char *arg_a[]){
  cudaSetDeviceFlags(cudaDeviceBlockingSync);
  if(arg_c<=1){
    listDevices();
    fprintf(stderr, "\nUse: %s [device] (f2k muons)\n"
	    "     %s [str] [om] [num] [device] (flasher)\n", arg_a[0], arg_a[0]);
  }
  else if(arg_c<=2){
    int device=0;
    if(arg_c>1) device=atoi(arg_a[1]);
    checkError(cudaSetDevice(device));
    checkError(cudaGetDeviceProperties(&prop, device));
    fprintf(stderr, "Processing f2k muons from stdin on device %d\n", device);
    f2k();
  }
  else{
    int str=0, dom=0, device=0, itr=0;
    unsigned long long num=1000000ULL;

    if(arg_c>1) str=atoi(arg_a[1]);
    if(arg_c>2) dom=atoi(arg_a[2]);
    if(arg_c>3){
      num=(unsigned long long) atof(arg_a[3]);
      char * sub = strchr(arg_a[3], '*');
      if(sub!=NULL) itr=(int) atof(++sub);
    }
    if(arg_c>4) device=atoi(arg_a[4]);
    checkError(cudaSetDevice(device));
    checkError(cudaGetDeviceProperties(&prop, device));
    fprintf(stderr, "Running flasher simulation on device %d\n", device);
    flasher(str, dom, num, itr);
  }

  fprintf(stderr, "\nDevice time: %2.1f [ms]\n", deviceTime);
  checkError(cudaThreadExit());
}