// -*-c++-*-
#include <curand_kernel.h>

__device__ unsigned int float_to_sortable_int(float f)
{
  return __float_as_int(f);
  //int i = __float_as_int(f);
  //unsigned int mask = -(int)(i >> 31) | 0x80000000;
  //return i ^ mask;
}

__device__ float sortable_int_to_float(unsigned int i)
{
  return __int_as_float(i);
  //unsigned int mask = ((i >> 31) - 1) | 0x80000000;
  //return __int_as_float(i ^ mask);
}

extern "C"
{

__global__ void reset_earliest_time_int(float maxtime, int ntime_ints, unsigned int *time_ints)
{
	int id = threadIdx.x + blockDim.x * blockIdx.x;
	if (id < ntime_ints) {
		unsigned int maxtime_int = float_to_sortable_int(maxtime);
		time_ints[id] = maxtime_int;
	}
}

__global__ void run_daq(curandState *s, unsigned int detection_state,
			float time_rms,
			int first_photon,
			int nphotons, float *photon_times,
			unsigned int *photon_histories,
			int *last_hit_triangles, int *solid_map,
			int nsolids, unsigned int *earliest_time_int,
			unsigned int *channel_q,
			unsigned int *channel_histories)
{

	int id = threadIdx.x + blockDim.x * blockIdx.x;

	if (id < nphotons)
	{
		curandState rng = s[id];
		int photon_id = id + first_photon;
		int triangle_id = last_hit_triangles[photon_id];
		
		if (triangle_id > -1)
		{
			int solid_id = solid_map[triangle_id];
			unsigned int history = photon_histories[photon_id];

			if (solid_id < nsolids && (history & detection_state)) 
			{
			  float time = photon_times[photon_id] + curand_normal(&rng) * time_rms;
			  unsigned int time_int = float_to_sortable_int(time);
			  
			  atomicMin(earliest_time_int + solid_id, time_int);
			  atomicAdd(channel_q + solid_id, 1);
			  atomicOr(channel_histories + solid_id, history);
			}

		}

		s[id] = rng;

	}

}

__global__ void convert_sortable_int_to_float(int n, 
					      unsigned int *sortable_ints, 
					      float *float_output)
{
	int id = threadIdx.x + blockDim.x * blockIdx.x;
	
	if (id < n)
		float_output[id] = sortable_int_to_float(sortable_ints[id]);
}

} // extern "C"