lagrangian

//*****************************************************************//
//** Kernel To Move Particles On the GPU independent of the grid **//
//*****************************************************************//
__global__ void Kernel_MoveParticles(dim *part_pos, dim *part_vel, dim *part_acc, int num_particles , double *timestep, int i_sp, bool glo_ESolver, int glostep, bool *dev_part_CEX, int CollInterval)
{
   int blockId = blockIdx.x + blockIdx.y*gridDim.x;
   int tid = blockId*(blockDim.x*blockDim.y) + threadIdx.y*blockDim.x + threadIdx.x;
   if(tid<num_particles)
   {
      if(glo_ESolver==1)
      {
         if(i_sp==0)
         {
            //**  for neutrals                                        **//
            //**  the acc here is typiclly 1, but after CEX           **//
            //**  acc is weight because of the speed of fast neutrals **//
//            part_pos[tid].x[0] += part_vel[tid].x[0]*timestep[i_sp]*part_acc[tid].x[0];
//            part_pos[tid].x[1] += part_vel[tid].x[1]*timestep[i_sp]*part_acc[tid].x[1];
//            part_pos[tid].x[2] += part_vel[tid].x[2]*timestep[i_sp]*part_acc[tid].x[2];
         }
         if(i_sp>0)
         {
            //** for ions and electrons                                           **//
            //** acc is initially 0, and will be calculated based on E = grad.Phi **//
            //** Using the Leap-frog Scheme                                       **//
             if(glostep==1)
             {
               dim dummvel;
//               dummvel.x[0] = part_vel[tid].x[0] - 0.5 * part_acc[tid].x[0] * timestep[i_sp];
//               dummvel.x[1] = part_vel[tid].x[1] - 0.5 * part_acc[tid].x[1] * timestep[i_sp];
//               dummvel.x[2] = part_vel[tid].x[2] - 0.5 * part_acc[tid].x[2] * timestep[i_sp];
//               part_vel[tid].x[0] = dummvel.x[0];
//               part_vel[tid].x[1] = dummvel.x[1];
//               part_vel[tid].x[2] = dummvel.x[2];
             }


             if(i_sp == 1)
             {
                if(dev_part_CEX[tid] == 1)
                {
//                 part_vel[tid].x[0] += part_acc[tid].x[0]*timestep[i_sp];//*CollInterval;
//                 part_vel[tid].x[1] += part_acc[tid].x[1]*timestep[i_sp];//*CollInterval;
//                 part_vel[tid].x[2] += part_acc[tid].x[2]*timestep[i_sp];//*CollInterval;
//                 part_pos[tid].x[0] += part_vel[tid].x[0]*timestep[i_sp];//*CollInterval;
//                 part_pos[tid].x[1] += part_vel[tid].x[1]*timestep[i_sp];//*CollInterval;
//                 part_pos[tid].x[2] += part_vel[tid].x[2]*timestep[i_sp];//*CollInterval;
                }
                else
                {
//                 part_vel[tid].x[0] += part_acc[tid].x[0]*timestep[i_sp];
//                 part_vel[tid].x[1] += part_acc[tid].x[1]*timestep[i_sp];
//                 part_vel[tid].x[2] += part_acc[tid].x[2]*timestep[i_sp];
//                 part_pos[tid].x[0] += part_vel[tid].x[0]*timestep[i_sp];
//                 part_pos[tid].x[1] += part_vel[tid].x[1]*timestep[i_sp];
//                 part_pos[tid].x[2] += part_vel[tid].x[2]*timestep[i_sp];
                }
             }
             else
             {

//              part_vel[tid].x[0] += part_acc[tid].x[0]*timestep[i_sp];
//              part_vel[tid].x[1] += part_acc[tid].x[1]*timestep[i_sp];
//              part_vel[tid].x[2] += part_acc[tid].x[2]*timestep[i_sp];
//              part_pos[tid].x[0] += part_vel[tid].x[0]*timestep[i_sp];
//              part_pos[tid].x[1] += part_vel[tid].x[1]*timestep[i_sp];
//              part_pos[tid].x[2] += part_vel[tid].x[2]*timestep[i_sp];
             }



         }
      }

      if(glo_ESolver==0)
      {
//         part_pos[tid].x[0] += part_vel[tid].x[0]*timestep[i_sp];
//         part_pos[tid].x[1] += part_vel[tid].x[1]*timestep[i_sp];
//         part_pos[tid].x[2] += part_vel[tid].x[2]*timestep[i_sp];
      }
   }
}