double Vector3D::get_norm() const {
	return eval_norm(_x, _y, _z);
} //returns the magnitude of a vector
int main(int argc, char **argv) {
  parse_cmd(argc, argv);

  cudaSetDevice(device);

  ffloat T = host_omega>0?(2*PI/host_omega):0; // period of external a/c emf
  if( display == 9 ) {
    t_max = t_start + 101*T;
    init_strobe_array();
  } else {
    t_max = t_start + T;
  }
  if( quiet == 0 ) { printf("# t_max = %0.20f kernel=%d\n", t_max, BLTZM_KERNEL); }

  // we will allocate enough memory to accommodate range
  // from -PhiY_max_range to PhiY_max_range, but will use only
  // part of it from PhiYmin to PhiYmax
  ffloat PhiY_max_range = fabs(PhiYmin);
  if( PhiY_max_range < fabs(PhiYmax) ) {
    PhiY_max_range = fabs(PhiYmax);
  }

  //host_dPhi = PhiY_max_range/host_M;
  host_dPhi = (PhiYmax-PhiYmin)/host_M;

  NSIZE = host_N+1;
  //MSIZE = 2*host_M+3;
  MSIZE = host_M+3;
  PADDED_MSIZE = (MSIZE*sizeof(ffloat))%128==0?MSIZE:((((MSIZE*sizeof(ffloat))/128)*128+128)/sizeof(ffloat));
  printf("PADDED MEMORY FROM %d ELEMENTS PER ROW TO %d\n", MSIZE, (int)PADDED_MSIZE);

  MP1 = host_M+1; // 

  SIZE_2D = NSIZE*PADDED_MSIZE;
  const int SIZE_2Df = SIZE_2D*sizeof(ffloat);

  host_TMSIZE=host_M+1;

  host_nu = 1+host_dt/2;
  host_nu2 = host_nu * host_nu;
  host_nu_tilde = 1-host_dt/2;
  host_bdt = host_B*host_dt/(4*host_dPhi);

  load_data();

  // create a0 and populate it with f0
  ffloat *host_a0; host_a0 = (ffloat *)calloc(SIZE_2D, sizeof(ffloat));
  for( int n=0; n<host_N+1; n++ ) {
    ffloat a = gsl_sf_bessel_In(n, host_mu)*(n==0?0.5:1)/(PI*gsl_sf_bessel_In(0, host_mu))*sqrt(host_mu/(2*PI*host_alpha));
    for( int m = 0; m < host_M+3; m++ ) {
      nm(host_a0, n, m) = a*expl(-host_mu*pow(phi_y(m),2)/2);
    }
  }

  // create device_a0 and transfer data from host_a0 to device_a0
  ffloat *a0;
  HANDLE_ERROR(cudaMalloc((void **)&a0, SIZE_2Df));
  HANDLE_ERROR(cudaMemcpy(a0, host_a0, SIZE_2Df, cudaMemcpyHostToDevice));

  // create a and b 2D vectors, four of each. one for current,
  // another for next pointer on main and shifted grids
  ffloat *host_a = (ffloat *)calloc(SIZE_2D, sizeof(ffloat));
  ffloat *host_b = (ffloat *)calloc(SIZE_2D, sizeof(ffloat));

  ffloat *a[4];
  ffloat *b[4];
  for( int i = 0; i < 4; i++ ) {
    HANDLE_ERROR(cudaMalloc((void **)&a[i], SIZE_2Df));
    HANDLE_ERROR(cudaMalloc((void **)&b[i], SIZE_2Df));

    // zero vector b[i]
    HANDLE_ERROR(cudaMemset((void *)a[i], 0, SIZE_2Df));
    HANDLE_ERROR(cudaMemset((void *)b[i], 0, SIZE_2Df));
  }

  int current = 0; int next = 1;
  int current_hs = 2; int next_hs = 3; // 'hs' - half step

  // init vectors a[0] and a[2]
  HANDLE_ERROR(cudaMemcpy(a[current], host_a0, SIZE_2Df,
                          cudaMemcpyHostToDevice));

  int blocks = (host_M+3)/TH_PER_BLOCK;

  // tiptow to the first half step
  ffloat *host_a_hs = (ffloat *)calloc(SIZE_2D, sizeof(ffloat));
  ffloat *host_b_hs = (ffloat *)calloc(SIZE_2D, sizeof(ffloat));
  ffloat cos_omega_t = 1; // cos(host_omega*t); for t = 0
  ffloat cos_omega_t_plus_dt = cos(host_omega*(host_dt));
  step_on_grid(blocks, a0, a[current], b[current], a[current_hs], b[current_hs],
               a[current], b[current], 0, 0,
               cos_omega_t, cos_omega_t_plus_dt);
  /*
  // temporary solution // FIX ME!!!
  memcpy(host_a_hs, host_a, SIZE_2D*sizeof(ffloat));
  HANDLE_ERROR(cudaMemcpy(a[current_hs], host_a_hs,
                          SIZE_2Df, cudaMemcpyHostToDevice));
  HANDLE_ERROR(cudaMemcpy(b[current_hs], host_b_hs,
                          SIZE_2Df, cudaMemcpyHostToDevice));
  */

  // used for file names when generated data for making animation
  char *file_name_buf = (char *)calloc(128, sizeof(char));

  char buf[16384]; // output buffer for writing frame data when display==77

  int step = 0;
  ffloat frame_time = 0; int frame_number = 1;

  ffloat *host_av_data; host_av_data = (ffloat *)calloc(5, sizeof(ffloat));
  ffloat *av_data;
  HANDLE_ERROR(cudaMalloc((void **)&av_data, 6*sizeof(ffloat)));
  HANDLE_ERROR(cudaMemset((void *)av_data, 0, 6*sizeof(ffloat)));

  float t_hs = 0;

  ffloat t0 = 0;
  ffloat t = t0;
  ffloat timeout = -999;

  ffloat last_tT_reminder = 0;

  for(;;) {
    //read_from
    int ccc = 0;
    for( t = t0; t < t_max; t += host_dt ) {
      /// XXX
      //ccc++;
      //if( ccc == 51 ) { break; }

      t_hs = t + host_dt/2;
      cos_omega_t = cos(host_omega*t);
      cos_omega_t_plus_dt = cos(host_omega*(t+host_dt));
      step_on_grid(blocks, a0, a[current], b[current], a[next], b[next], a[current_hs],
                   b[current_hs], t, t_hs,
                   cos_omega_t, cos_omega_t_plus_dt);

      cudaThreadSynchronize();

      cos_omega_t = cos(host_omega*t_hs);
      cos_omega_t_plus_dt = cos(host_omega*(t_hs+host_dt));
      step_on_half_grid(blocks, a0, a[current], b[current], a[next], b[next], a[current_hs],
                        b[current_hs], a[next_hs], b[next_hs], t, t_hs,
                        cos_omega_t, cos_omega_t_plus_dt);

      /*
      if( t >= 0 ) { 
	HANDLE_ERROR(cudaMemcpy(host_a, a[current], SIZE_2Df, cudaMemcpyDeviceToHost));
	HANDLE_ERROR(cudaMemcpy(host_b, b[current], SIZE_2Df, cudaMemcpyDeviceToHost));
          sprintf(file_name_buf, "strobe.data");
          FILE *frame_file_stream = fopen(file_name_buf, "w");
          setvbuf(frame_file_stream, buf, _IOFBF, sizeof(buf));
          printf("\nWriting strobe %s\n", file_name_buf);
          print_2d_strobe(frame_file_stream, MSIZE, host_a0, host_a, host_b, host_alpha, t);
          fclose(frame_file_stream);
          frame_time = 0;

	break; } /// XXX REMOVE ME
      */

      if( host_E_omega > 0 && display == 77 && frame_time >= 0.01) {
        // we need to perform averaging of v_dr, m_x and A
        av(blocks, a[next], b[next], av_data, t);
        HANDLE_ERROR(cudaMemcpy(host_a, a[current], SIZE_2Df, cudaMemcpyDeviceToHost));
        HANDLE_ERROR(cudaMemcpy(host_b, b[current], SIZE_2Df, cudaMemcpyDeviceToHost));
        HANDLE_ERROR(cudaMemcpy(host_av_data, av_data, 6*sizeof(ffloat), cudaMemcpyDeviceToHost));
        ffloat norm = eval_norm(host_a, host_alpha, MSIZE);
        print_time_evolution_of_parameters(out, norm, host_a, host_b, MSIZE,
                                           host_mu, host_alpha, host_E_dc, host_E_omega, host_omega,
                                           host_av_data, t);
        frame_time = 0;
      }

      if( host_E_omega > 0 && display != 7 && display != 77 && display != 8 && t >= t_start ) {
        // we need to perform averaging of v_dr, m_x and A
        av(blocks, a[next], b[next], av_data, t);
      }

      if( current    == 0 ) {    current = 1;    next = 0; } else { current = 0; next = 1; }
      if( current_hs == 2 ) { current_hs = 3; next_hs = 2; } else { current_hs = 2; next_hs = 3; }

      //if( display == 9 && t >= t_start ) {
      //  ffloat tT = t/T;
      //  printf("t=%0.12f %0.12f %0.12f\n", t, , T);
      //}

      if( display == 9 && t >= t_start ) { // XXX PUT ME BACK
        ffloat tT = t/T;
        ffloat tT_reminder = tT-((int)tT);
        if( tT_reminder < last_tT_reminder ) { 
          HANDLE_ERROR(cudaMemcpy(host_a, a[current], SIZE_2Df, cudaMemcpyDeviceToHost));
          HANDLE_ERROR(cudaMemcpy(host_b, b[current], SIZE_2Df, cudaMemcpyDeviceToHost));
          sprintf(file_name_buf, "strobe%08d.data", frame_number++);
          FILE *frame_file_stream = fopen(file_name_buf, "w");
          setvbuf(frame_file_stream, buf, _IOFBF, sizeof(buf));
          printf("\nWriting strobe %s\n", file_name_buf);
          print_2d_strobe(frame_file_stream, MSIZE, host_a0, host_a, host_b, host_alpha, t);
          fclose(frame_file_stream);
          frame_time = 0;
	}
        last_tT_reminder = tT_reminder;
      }

      if( display == 7 && frame_time >= 0.01 && t > frame_start ) { // we are making movie
        HANDLE_ERROR(cudaMemcpy(host_a, a[current], SIZE_2Df, cudaMemcpyDeviceToHost));
        HANDLE_ERROR(cudaMemcpy(host_b, b[current], SIZE_2Df, cudaMemcpyDeviceToHost));
        sprintf(file_name_buf, "frame%08d.data", frame_number++);
        FILE *frame_file_stream = fopen(file_name_buf, "w");
        setvbuf(frame_file_stream, buf, _IOFBF, sizeof(buf));
        printf("\nWriting frame %s\n", file_name_buf);
        print_2d_data(frame_file_stream, MSIZE, host_a0, host_a, host_b, host_alpha, t);
        fclose(frame_file_stream);
        frame_time=0;
      }

      if( out != stdout && display != 7 ) {
        step++;
        if( step == 300 ) {
          printf("\rt=%0.9f %0.2f%%", t, t/t_max*100);
          fflush(stdout);
          step = 0;
        }
      }
      frame_time += host_dt;

      if( display == 9 && t <= t_start && frame_time >= T ) {
        frame_time == 0;
      }
    }

    HANDLE_ERROR(cudaMemcpy(host_a, a[current], SIZE_2Df, cudaMemcpyDeviceToHost));
    HANDLE_ERROR(cudaMemcpy(host_b, b[current], SIZE_2Df, cudaMemcpyDeviceToHost));
    HANDLE_ERROR(cudaMemcpy(host_av_data, av_data, 6*sizeof(ffloat), cudaMemcpyDeviceToHost));

    ffloat norm = 0;
    ffloat dphi_over_2 = host_dPhi/2.0;
    for( int m = 1; m < host_M+1; m++ ) {
      norm += (nm(host_a,0,m)+nm(host_a,0,m))*dphi_over_2;
    }
    norm *= 2*PI*sqrt(host_alpha);

    if( display == 3 ) {
      for( ffloat phi_x = -PI; phi_x < PI; phi_x += 0.01 ) {
        for( int m = 1; m < host_M; m++ ) {
          ffloat value = 0;
          ffloat value0 = 0;
          for( int n = 0; n < host_N+1; n++ ) {
            value  += nm(host_a,n,m)*cos(n*phi_x) + nm(host_b,n,m)*sin(n*phi_x);
            value0 += nm(host_a0,n,m)*cos(n*phi_x);
          }
          fprintf(out, "%0.5f %0.5f %0.20f %0.20f\n", phi_x, phi_y(m), value<0?0:value, value0<0?0:value0);
        }
      }
      fprintf(out, "# norm=%0.20f\n", norm);
      printf("# norm=%0.20f\n", norm);
      //if( out != stdout ) { fclose(out); }
      cuda_clean_up();
      return EXIT_SUCCESS;
    }

    if( display == 8 ) {
      // single shot image
      HANDLE_ERROR(cudaMemcpy(host_a, a[current], SIZE_2Df, cudaMemcpyDeviceToHost));
      HANDLE_ERROR(cudaMemcpy(host_b, b[current], SIZE_2Df, cudaMemcpyDeviceToHost));
      sprintf(file_name_buf, "frame.data");
      FILE *frame_file_stream = fopen(file_name_buf, "w");
      setvbuf(frame_file_stream, buf, _IOFBF, sizeof(buf));
      printf("\nWriting frame %s\n", file_name_buf);
      print_2d_data(frame_file_stream, MSIZE, host_a0, host_a, host_b, host_alpha, t);
      fclose(frame_file_stream);
      frame_time=0;
      return EXIT_SUCCESS;
    }

    if( display == 4 ) {
      if( quiet == 0 ) { printf("\n# norm=%0.20f\n", norm); }
      ffloat v_dr_inst = 0 ;
      ffloat v_y_inst = 0;
      ffloat m_over_m_x_inst = 0;
      for( int m = 1; m < host_M; m++ ) {
        v_dr_inst += nm(host_b,1,m)*host_dPhi;
        v_y_inst  += nm(host_a,0,m)*phi_y(m)*host_dPhi;
        m_over_m_x_inst += nm(host_a,1,m)*host_dPhi;
      }

      ffloat v_dr_multiplier = 2*gsl_sf_bessel_I0(host_mu)*PI*sqrt(host_alpha)/gsl_sf_bessel_In(1, host_mu);
      ffloat v_y_multiplier  = 4*PI*gsl_sf_bessel_I0(host_mu)/gsl_sf_bessel_In(1, host_mu);
      ffloat m_over_multiplier = PI*host_alpha*sqrt(host_alpha);
      v_dr_inst       *= v_dr_multiplier;
      v_y_inst        *= v_y_multiplier;
      m_over_m_x_inst *= m_over_multiplier;

      host_av_data[1] *= v_dr_multiplier;
      host_av_data[2] *= v_y_multiplier;
      host_av_data[3] *= m_over_multiplier;
      host_av_data[4] *= v_dr_multiplier;
      host_av_data[4] /= T;
      host_av_data[5] *= v_dr_multiplier;
      host_av_data[5] /= T;

      fprintf(out, "# display=%d E_dc=%0.20f E_omega=%0.20f omega=%0.20f mu=%0.20f alpha=%0.20f n-harmonics=%d PhiYmin=%0.20f PhiYmax=%0.20f B=%0.20f t-max=%0.20f dt=%0.20f g-grid=%d\n",
                      display,   host_E_dc,  host_E_omega,  host_omega,  host_mu,  host_alpha,  host_N,        PhiYmin,       PhiYmax,       host_B,  t_start,     host_dt,  host_M);
      fprintf(out, "#E_{dc}                \\tilde{E}_{\\omega}     \\tilde{\\omega}         mu                     v_{dr}/v_{p}         A(\\omega)              NORM     v_{y}/v_{p}    m/m_{x,k}   <v_{dr}/v_{p}>   <v_{y}/v_{p}>    <m/m_{x,k}>    Asin\n");
      fprintf(out, "%0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f %0.20f\n",
              host_E_dc, host_E_omega, host_omega, host_mu, v_dr_inst, host_av_data[4], norm, v_y_inst,
              m_over_m_x_inst, host_av_data[1], host_av_data[2], host_av_data[3], host_av_data[5]);
    }

    if( read_from == NULL ) { break; }

    // scan for new parameters
    timeout = scan_for_new_parameters();
    if( timeout < -900 ) { break; } // user entered 'exit'
    t_start = t + timeout;
    t_max = t_start + T;
    t0 = t + host_dt;
    T=host_omega>0?(2*PI/host_omega):0;
    load_data(); // re-load data
    HANDLE_ERROR(cudaMemset((void *)av_data, 0, 6*sizeof(ffloat))); // clear averaging data
    if( quiet == 0 ) { printf("# t_max = %0.20f\n", t_max); }
  } // for(;;)

  if( out != NULL && out != stdout ) {
    fclose(out);
  }
  cuda_clean_up();
  return EXIT_SUCCESS;
} // end of main(...)
double HcurlOrthoHP::calc_error_n(int n, ...)
{
  int i, j, k;

  if (n != num) error("Wrong number of solutions.");

  // obtain solutions and bilinear forms
  va_list ap;
  va_start(ap, n);
  for (i = 0; i < n; i++) {
    sln[i] = va_arg(ap, Solution*);
    sln[i]->set_quad_2d(&g_quad_2d_std);
  }
  for (i = 0; i < n; i++) {
    rsln[i] = va_arg(ap, Solution*);
    rsln[i]->set_quad_2d(&g_quad_2d_std);
  }
  va_end(ap);

  // prepare multi-mesh traversal and error arrays
  AUTOLA_OR(Mesh*, meshes, 2*num);
  AUTOLA_OR(Transformable*, tr, 2*num);
  Traverse trav;
  nact = 0;
  for (i = 0; i < num; i++)
  {
    meshes[i] = sln[i]->get_mesh();
    meshes[i+num] = rsln[i]->get_mesh();
    tr[i] = sln[i];
    tr[i+num] = rsln[i];

    nact += sln[i]->get_mesh()->get_num_active_elements();

    int max = meshes[i]->get_max_element_id();
    if (errors[i] != NULL) delete [] errors[i];
    errors[i] = new double[max];
    memset(errors[i], 0, sizeof(double) * max);
  }

  double total_norm = 0.0;
  AUTOLA_OR(double, norms, num);
  memset(norms, 0, num*sizeof(double));
  double total_error = 0.0;
  if (esort != NULL) delete [] esort;
  esort = new int2[nact];

  Element** ee;
  trav.begin(2*num, meshes, tr);
  while ((ee = trav.get_next_state(NULL, NULL)) != NULL)
  {
    for (i = 0; i < num; i++)
    {
      RefMap* rmi = sln[i]->get_refmap();
      RefMap* rrmi = rsln[i]->get_refmap();
      for (j = 0; j < num; j++)
      {
        RefMap* rmj = sln[j]->get_refmap();
        RefMap* rrmj = rsln[j]->get_refmap();
        double e, t;
        if (form[i][j] != NULL)
        {
          #ifndef COMPLEX
          e = fabs(eval_error(form[i][j], ord[i][j], sln[i], sln[j], rsln[i], rsln[j], rmi, rmj, rrmi, rrmj));
          t = fabs(eval_norm(form[i][j], ord[i][j], rsln[i], rsln[j], rrmi, rrmj));
          #else
          e = std::abs(eval_error(form[i][j], ord[i][j], sln[i], sln[j], rsln[i], rsln[j], rmi, rmj, rrmi, rrmj));
          t = std::abs(eval_norm(form[i][j], ord[i][j], rsln[i], rsln[j], rrmi, rrmj));
          #endif

          norms[i] += t;
          total_norm  += t;
          total_error += e;
          errors[i][ee[i]->id] += e;
        }
      }
    }
  }
  trav.finish();

  Element* e;
  k = 0;
  for (i = 0; i < num; i++)
    for_all_active_elements(e, meshes[i]) {
      esort[k][0] = e->id;
      esort[k++][1] = i;
      errors[i][e->id] /= norms[i];
    }