#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <vector>
#include <fstream>

#include "SearchKernel.cuh"
using namespace std;

#ifndef MAX
#define MAX(a,b) (a > b ? a : b)
#endif

#define CUDA_CHECK_RETURN(value) {											\
	cudaError_t _m_cudaStat = value;										\
	if (_m_cudaStat != cudaSuccess) {										\
		fprintf(stderr, "Error %s at line %d in file %s\n",					\
				cudaGetErrorString(_m_cudaStat), __LINE__, __FILE__);		\
		exit(1);															\
	} }


typedef struct __align__(16)  {
	int startPosition;
	int hitsNum;
	double z;
} sensorInfo;

struct track {
	double m_x0;
	double m_tx;
	double m_y0;
	double m_ty;

	double m_s0;
	double m_sx;
	double m_sz;
	double m_sxz;
	double m_sz2;

	double m_u0;
	double m_uy;
	double m_uz;
	double m_uyz;
	double m_uz2;

	int internalId;
	int trackHitsNum; // hmm
	int firstHit; 	 //	hmm

	bool m_backward;
};

__device__ __constant__ double 	m_maxXSlope			= 0.400;
__device__ __constant__ double 	m_maxYSlope			= 0.300;
__device__ __constant__ double 	m_maxZForRBeamCut	= 200.0;
__device__ __constant__ double 	m_maxR2Beam			= 1.0 ;
__device__ __constant__ int 	m_maxMissed			= 4;
__device__ __constant__ double 	m_extraTol			= 0.150;
__device__ __constant__ double 	m_maxChi2ToAdd		= 100.0;
__device__ __constant__ double 	m_maxChi2SameSensor	= 16.0;
__device__ __constant__ double  m_maxChi2Short		= 6.0 ;
__device__ __constant__ double  m_maxChi2PerHit		= 16.0;
__device__ __constant__ int 	m_sensNum			= 48;


__device__ double zBeam(track *tr) {
	return -( tr->m_x0 * tr->m_tx + tr->m_y0 * tr->m_ty ) / ( tr->m_tx * tr->m_tx + tr->m_ty * tr->m_ty );
}

__device__ double r2AtZ( double z , track *tr) {
    double xx = tr->m_x0 + z * tr->m_tx;
    double yy = tr->m_y0 + z * tr->m_ty;
    return xx*xx + yy * yy;
 }

__device__ void solve (track *tr) {
	double den = ( tr->m_sz2 * tr->m_s0 - tr->m_sz * tr->m_sz );
	if ( fabs(den) < 10e-10 ) den = 1.;
	tr->m_tx     = ( tr->m_sxz * tr->m_s0  - tr->m_sx  * tr->m_sz ) / den;
	tr->m_x0     = ( tr->m_sx  * tr->m_sz2 - tr->m_sxz * tr->m_sz ) / den;

	den = ( tr->m_uz2 * tr->m_u0 - tr->m_uz * tr->m_uz );
	if ( fabs(den) < 10e-10 ) den = 1.;
	tr->m_ty     = ( tr->m_uyz * tr->m_u0  - tr->m_uy  * tr->m_uz ) / den;
	tr->m_y0     = ( tr->m_uy  * tr->m_uz2 - tr->m_uyz * tr->m_uz ) / den;
}

__device__ void addHit ( track *tr, int offset, int *trackIds, double *inX, double *inY, double *inZ, double *inW) {

	trackIds[offset] = tr->internalId;
	tr->trackHitsNum++;

	double z = inZ[offset];
	double x = inX[offset];
	double w = inW[offset];

	tr->m_s0  += w;
	tr->m_sx  += w * x;
	tr->m_sz  += w * z;
	tr->m_sxz += w * x * z;
	tr->m_sz2 += w * z * z;

	double y = inY[offset];

	tr->m_u0  += w;
	tr->m_uy  += w * y;
	tr->m_uz  += w * z;
	tr->m_uyz += w * y * z;
	tr->m_uz2 += w * z * z;

	if( tr->trackHitsNum > 1 ) solve(tr);
}

__device__ void setTrack(track *tr, int hit0offset, int hit1offset, int *trackIds, double *inX, double *inY, double *inZ, double *inW){
	tr->m_backward = false;
	tr->trackHitsNum = 0;

	trackIds[hit0offset] = tr->internalId;
	tr->trackHitsNum++;

	double z = inZ[hit0offset];
	double x = inX[hit0offset];
	double w = inW[hit0offset];

	tr->m_s0  = w;
	tr->m_sx  = w * x;
	tr->m_sz  = w * z;
	tr->m_sxz = w * x * z;
	tr->m_sz2 = w * z * z;

	double y = inY[hit0offset];

	tr->m_u0  = w;
	tr->m_uy  = w * y;
	tr->m_uz  = w * z;
	tr->m_uyz = w * y * z;
	tr->m_uz2 = w * z * z;

	addHit (tr, hit1offset,  trackIds, inX, inY, inZ, inW);
}

__device__ inline double chi2Hit( double x, double y, double hitX, double hitY, double hitW){
	double dx = x - hitX;
	double dy = y - hitY;
	return dx * dx * (hitW) + dy * dy * (hitW);
}
__device__ inline double xAtHit(track *tr, double z )
{
	return tr->m_x0 + tr->m_tx * z;
}

__device__ inline double yAtHit( track *tr, double z  )
{
	return tr->m_y0 + tr->m_ty * z;
}

__device__ inline double chi2Track(track *tr, double hitX, double hitY, double hitZ, double hitW)
{
	return chi2Hit( xAtHit( tr, hitZ ), yAtHit(tr, hitZ ), hitX, hitY, hitW);
}

__device__ inline double chi2(track *tr, double *hitX, double *hitY, double *hitZ, double *hitW, int usedHits[])
{
	double ch = 0.0;
	int nDoF  = -4;
	for (int i =0 ; i<m_sensNum; i++){ // or just... while(i!=m_sensNum && usedHits[i] != -1) ,
		int hitNumber = usedHits[i];
		if(hitNumber== -1){
			break;
		}
		ch += chi2Track(tr, hitX[hitNumber], hitY[hitNumber], hitZ[hitNumber], hitW[hitNumber]);
		nDoF += 2;
	}
	return ch/nDoF;
}

__device__ inline bool addHitsOnSensor( double sensZ, int sensStartPost, int sensHitsNum, double xTol, double maxChi2, track *tr,
								 	    int *tracksIds, double *inX, double *inY, double *inZ, double *inW, int threadId,
								 	    int threadsNumber,int *usedHit ) {
	if (sensHitsNum == 0 ) return false;
	double xGuess = (tr->m_x0 + tr->m_tx * sensZ) - xTol - 1;
	int lastHit = sensStartPost + sensHitsNum -1;
	if((inX[(lastHit*threadsNumber)+threadId]) < xGuess) return false;
	int startHit = sensStartPost; // może lepiej bezpośrednio się do tego odnosić
	unsigned int step = sensHitsNum;
	while ( step > 2 ) {
		step = step/2;
		if (inX[((startHit+step)*threadsNumber)+threadId] < xGuess) startHit +=step;
	}
	bool added = false;
	int tmpOffset = 0;
	for(int iH = startHit; iH<=lastHit ;++iH){ ////
		tmpOffset = (iH*threadsNumber)+threadId;
		double xPred = (tr->m_x0 + tr->m_tx * inZ[tmpOffset]);
		if ( inX[tmpOffset] + xTol < xPred ) continue;
		if ( inX[tmpOffset] - xTol > xPred ) break;
		if ( chi2Track(tr, inX[tmpOffset], inY[tmpOffset], inZ[tmpOffset], inW[tmpOffset]) < maxChi2 ) {
			addHit(tr, tmpOffset,tracksIds, inX, inY, inZ, inW );
			*usedHit = tmpOffset;
			added = true;
		}
	}
	return added;
}

__device__ inline void removeHit(track *tr, int worstHitOffset, double *inX, double *inY, double *inZ, double *inW ){
	tr->trackHitsNum--;

	double z = inZ[worstHitOffset];
	double w = inW[worstHitOffset];
	double x = inX[worstHitOffset];

	tr->m_s0  -= w;
	tr->m_sx  -= w * x;
	tr->m_sz  -= w * z;
	tr->m_sxz -= w * x * z;
	tr->m_sz2 -= w * z * z;

	double y = inY[worstHitOffset];

	tr->m_u0  -= w;
	tr->m_uy  -= w * y;
	tr->m_uz  -= w * z;
	tr->m_uyz -= w * y * z;
	tr->m_uz2 -= w * z * z;

	if( tr->trackHitsNum > 1 ) solve(tr);
}

__device__ inline void removeWorstHit(track* tr, double maxChi2, int usedHits[], int *tracksIds, double *inX, double *inY, double *inZ,
									  double *inW, int *isUsed, int threadsNumber, int threadId)
{
	double topChi2 = 1.e9;
	while( topChi2 > maxChi2 ) {
	    topChi2 = 0.0;
	    int offset = 0;
	    for (int i =0 ; i<m_sensNum; i++){
	    	offset = usedHits[i];
	    	if(offset == -1){
	    		break;
	    	}
	    	double myChi2 = chi2Track(tr, inX[offset], inY[offset], inZ[offset], inW[offset] );
	    	if ( myChi2 > topChi2 ) {
	    		topChi2 = myChi2;
	    		offset = (usedHits[i]*threadsNumber)+threadId;
	    	}
	    }
	    if ( topChi2 > maxChi2 ) {
	      tracksIds[offset] = -1;
	      isUsed[offset] = 0;
	      removeHit(tr, offset, inX, inY, inZ, inW);
	    }
	}
}

__device__ inline bool all3SensorsAreDifferent(int usedHits[], int *sensorsIds) {
    if ( sensorsIds[usedHits[0]] == sensorsIds[usedHits[1]]) return false;
    if ( sensorsIds[usedHits[0]] == sensorsIds[usedHits[2]]) return false;
    if ( sensorsIds[usedHits[1]] == sensorsIds[usedHits[2]]) return false;
    return true;
}

__device__ inline int nbUnused(int usedHits[],int *isUsed) {
	int nn = 0;
    for (int i=0; i<m_sensNum; i++){
    	int hitNumber = usedHits[i];
    	if(hitNumber== -1){
    		break;
    	}
    	if ( ! isUsed[hitNumber] ) ++nn;
    }
    return nn;
}

/*
 * topHitsNum - number of hits in the event with biggest number of hits. Other events are filled with dummy hits to have the same number of hits.
 */
__global__ void searchByPair(void *data, void* resultsKernel, int topHitsNum){

	int threadId = blockDim.x * blockIdx.x + threadIdx.x;
	track m_track = {0};			// move to shared memory
	int trackId = 0;
	int nextHitNb = 0;
	int usedHits[48] = {-1}; //numbers of used hits

	__shared__ int 		*inHitsNum;
	__shared__ int 		*inSensInfoStartPos;
	__shared__ int 		*inSensInfoHitsNum;
	__shared__ double 	*inSensInfoZ;
	__shared__ int 		*hitsIds;
	__shared__ int 		*tracksIds;
	__shared__ int 		*sensorsIds;
	__shared__ int 		*isUsed;
	__shared__ double 	*inX;
	__shared__ double 	*inY;
	__shared__ double 	*inZ;
	__shared__ double 	*inW;

	int eventsNumber = gridDim.x * blockDim.x;			//number of all the events passed to GPU
	int dataOffset =  eventsNumber * topHitsNum; 		// between different hits parameters

	if(threadId % blockDim.x == 0){	//only first thread from the block does that - maybe not the best idea ;)
		inHitsNum 				= 		(int*)data;
		inSensInfoStartPos 		= 		inHitsNum + eventsNumber;
		inSensInfoHitsNum		= 		inSensInfoStartPos + (eventsNumber * m_sensNum);
		inSensInfoZ 			= 		(double*)(inSensInfoHitsNum + (eventsNumber * m_sensNum));
		hitsIds 				= 		(int*)(inSensInfoZ + (eventsNumber * m_sensNum));
		tracksIds 				= 		hitsIds + dataOffset;
		sensorsIds				=		tracksIds + dataOffset;
		isUsed					=		sensorsIds + dataOffset;
		inX 					= 		(double*) (isUsed + dataOffset);
		inY 					= 		inX + dataOffset;
		inZ						=		inY + dataOffset;
		inW						=		inZ + dataOffset;
	}
	__syncthreads();


//	int hitsNum = inHitsNum[threadId];					//real number of hits for this event


	track *resTracks = (track*)(resultsKernel);			//tracks are stored as array of structs - should be stored as arrays of ints (coalesced access)
	sensorInfo sensor0;	//not the best solution
	sensorInfo sensor1;

	int lastSensor = m_sensNum-1;
	int firstSensor = 2;
	for ( int sens0 = lastSensor; firstSensor <= sens0; sens0 -= 1 ) {
		int sens1 = sens0 - 2;
		sensor0.startPosition = inSensInfoStartPos[threadId+(sens0*eventsNumber)];
		sensor0.hitsNum	= inSensInfoHitsNum[threadId+(sens0*eventsNumber)];
		sensor0.z = inSensInfoZ[threadId+(sens0*eventsNumber)];

		sensor1.startPosition = inSensInfoStartPos[threadId+(sens1*eventsNumber)];
		sensor1.hitsNum	= inSensInfoHitsNum[threadId+(sens1*eventsNumber)];
		sensor1.z = inSensInfoZ[threadId+(sens1*eventsNumber)];

		int hit0offset;
		int hit1offset;

		double dxMax = m_maxXSlope * fabs( sensor1.z - sensor0.z );
		double dyMax = m_maxYSlope * fabs( sensor1.z - sensor0.z );

		int first1 = sensor1.startPosition;
		for(int i0=sensor0.startPosition; i0<(sensor0.startPosition + sensor0.hitsNum );++i0){
			hit0offset = (i0*eventsNumber)+threadId;
			if(isUsed[hit0offset]){
				continue;
			}
			double x0 = inX[hit0offset];  //TODO: check if it goes to register of stays in local (global) memory
			double y0 = inY[hit0offset];
			double xMin = x0 - dxMax;
			double xMax = x0 + dxMax;
			for(int i1 = first1; i1<(sensor1.startPosition + sensor1.hitsNum ); i1++){
				memset(usedHits, -1, m_sensNum*sizeof(int));	//czy to na pewno tutaj ma być?
				hit1offset = (i1*eventsNumber)+threadId;
				double x1 = inX[hit1offset];
				if(x1<xMin){
					first1 = i1+1;
					continue;
				}
				if (x1 > xMax) break;
				if (isUsed[hit1offset]) continue;

				double y1  = inY[hit1offset];
				if ( fabs( y1 - y0 ) > dyMax ) continue;

				m_track.internalId = trackId;
				m_track.firstHit = nextHitNb;

				setTrack(&m_track, hit0offset, hit1offset, tracksIds , inX, inY, inZ, inW);
				usedHits[0] = hit0offset;   //!!!
				usedHits[1] = hit1offset;	//!!!

		        if ( sensor0.z < m_maxZForRBeamCut ) {
		          double z_beam  = zBeam(&m_track);
		          if ( z_beam > sensor0.z ) {
		            double r2Beam = r2AtZ( z_beam, &m_track );
		            if ( r2Beam > m_maxR2Beam )  continue;
		          }
		        }

		        int extraStep = 2;
		        int extraSens = sens1-extraStep;
		        int nbMissed = 0;
		        int extraHitId = 2;
		        double lastZ = sensor1.z;
		        //double lastZ = inSensInfoZ[threadId + (sens1*eventsNumber)];
		        while ( extraSens >= 0 ) {
		            double tol     =  m_extraTol;
		            double maxChi2 =  m_maxChi2ToAdd;
		            if ( inSensInfoZ[threadId + (extraSens*eventsNumber)] < lastZ - 100.0 ) {
		            	tol     = 2 * tol;
		                maxChi2 = 2 * maxChi2;
		            }
		            bool added = addHitsOnSensor(inSensInfoZ[threadId+extraSens],inSensInfoStartPos[threadId+extraSens],inSensInfoHitsNum[threadId+extraSens], tol, maxChi2,
		            							&m_track, tracksIds, inX, inY, inZ,inW, threadId, eventsNumber, &usedHits[extraHitId]);
		            if ( added ) {
		            	extraHitId++;
		            	nbMissed = 0;
		                lastZ = inSensInfoZ[threadId + extraSens];
		            } else {
		            	nbMissed += extraStep;
		                extraStep = 1;
		            }
		            if ( m_maxMissed < nbMissed ) break;
		            	extraSens -= extraStep;
		         }

		        //== Try upstream if almost forward tracks
		        if ( sensor0.z > m_maxZForRBeamCut ) {
		          extraStep = 1;
		          extraSens = sens0 + 3;
		          nbMissed = 2;
		          while ( extraSens <= lastSensor ) {
		        	  //extra = &sensors[extraSens];
		        	  int sensOffset = threadId+(extraSens*eventsNumber);
		        	  bool added = addHitsOnSensor(inSensInfoZ[sensOffset],inSensInfoStartPos[sensOffset],inSensInfoHitsNum[sensOffset], m_extraTol, m_maxChi2ToAdd,
		        	  		            		   &m_track, tracksIds, inX, inY, inZ,inW, threadId, eventsNumber, &usedHits[extraHitId]);
		            if ( added ) {
		              nbMissed = 0;
		              extraHitId++;
		            } else {
		              nbMissed += extraStep;
		            }
		            if ( m_maxMissed < nbMissed ) break;
		            extraSens += extraStep;
		          }
		        }
		        removeWorstHit(&m_track,m_maxChi2PerHit,usedHits, tracksIds, inX, inY, inZ, inW, isUsed,eventsNumber,threadId );
		        if ( m_track.trackHitsNum < 3 ) continue;
		        //== Add compatible hits in sens0 and sens1.
		        int tmpOffset;
		        if(i0 != (sensor0.startPosition + sensor0.hitsNum -1)){
		        	tmpOffset = ((i0+1)*eventsNumber)+threadId ;
		        	if ( chi2Track(&m_track, inX[tmpOffset], inY[tmpOffset],inZ[tmpOffset],inW[tmpOffset]) < m_maxChi2SameSensor) {
		        		++i0;
		                addHit(&m_track, hit0offset, tracksIds, inX, inY, inZ, inW );
		            }
		        }
		        if (i1 != (sensor1.startPosition + sensor0.hitsNum -1) ) {
		        	tmpOffset = ((i1+1)*eventsNumber)+threadId ;
		        	if ( chi2Track(&m_track, inX[tmpOffset], inY[tmpOffset],inZ[tmpOffset],inW[tmpOffset]) < m_maxChi2SameSensor) {
		        		++i1;
		        		addHit(&m_track, hit1offset, tracksIds, inX, inY, inZ, inW );
		            }
		        }
		        //== Final check: if only 3 hits, all should be unused and chi2 good.
		        if ( m_track.trackHitsNum == 3 ) {
		        	if ( !all3SensorsAreDifferent(usedHits, sensorsIds) ) {
		                continue;
		            }
		        	if(nbUnused(usedHits,isUsed) != 3){
		        		continue;
		            }
		            if(chi2(&m_track, inX, inY, inZ, inW, usedHits) > m_maxChi2Short){
		            	continue;
		            }
		        } else {
		        	if ( nbUnused(usedHits,isUsed) < .6 * m_track.trackHitsNum ) {
		        		continue;
		        	}
		        }
		        resTracks[trackId] = m_track; ///--- numer ścieżki

				trackId++;
				nextHitNb += m_track.trackHitsNum;

				if ( m_track.trackHitsNum > 3 ) {
					for (int i =0 ; i<48; i++){
						int hitNumber = usedHits[i];
						if(hitNumber== -1){
							break;
					 	}
						isUsed[hitNumber] = 1;
				    }
					break;
				}
			} //i1
		} //i0
	} //sensor0
//	resStats[0] = trackId;
}

inline int _ConvertSMVer2Cores(int major, int minor)
{
    // defines for GPU Architecture types
    typedef struct
    {
        int SM; // 0xMm, M = SM Major version, and m = SM minor version
        int Cores;
    } sSMtoCores;

    sSMtoCores nGpuArchCoresPerSM[] =
    {
        { 0x10,  8 }, // Tesla Generation (SM 1.0) G80 class
        { 0x11,  8 }, // Tesla Generation (SM 1.1) G8x class
        { 0x12,  8 }, // Tesla Generation (SM 1.2) G9x class
        { 0x13,  8 }, // Tesla Generation (SM 1.3) GT200 class
        { 0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class
        { 0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class
        { 0x30, 192}, // Kepler Generation (SM 3.0) GK10x class
        { 0x35, 192}, // Kepler Generation (SM 3.5) GK11x class
    };

    int index = 0;
    while (nGpuArchCoresPerSM[index].SM != -1)
    {
        if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor))
        {
            return nGpuArchCoresPerSM[index].Cores;
        }
        index++;
    }
    return nGpuArchCoresPerSM[7].Cores;
}

inline int getBestDevice()
{
    int current_device     = 0, sm_per_multiproc  = 0;
    int max_compute_perf   = 0, max_perf_device   = 0;
    int device_count       = 0, best_SM_arch      = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceCount(&device_count);

    // Find the best major SM Architecture GPU device
    while (current_device < device_count)
    {
        cudaGetDeviceProperties(&deviceProp, current_device);
        if (deviceProp.computeMode != cudaComputeModeProhibited)
        {
            if (deviceProp.major > 0 && deviceProp.major < 9999)
            {
                best_SM_arch = MAX(best_SM_arch, deviceProp.major);
            }
        }
        current_device++;
    }
    // Find the best CUDA capable GPU device
    current_device = 0;
    while (current_device < device_count)
    {
        cudaGetDeviceProperties(&deviceProp, current_device);
        // If this GPU is not running on Compute Mode prohibited, then we can add it to the list
        if (deviceProp.computeMode != cudaComputeModeProhibited)
        {
            if (deviceProp.major == 9999 && deviceProp.minor == 9999)
            {
                sm_per_multiproc = 1;
            }
            else
            {
                sm_per_multiproc = _ConvertSMVer2Cores(deviceProp.major, deviceProp.minor);
            }

            int compute_perf  = deviceProp.multiProcessorCount * sm_per_multiproc * deviceProp.clockRate;

            if (compute_perf  > max_compute_perf)
            {
                // If we find GPU with SM major > 2, search only these
                if (best_SM_arch > 2)
                {
                    // If our device==dest_SM_arch, choose this, or else pass
                    if (deviceProp.major == best_SM_arch)
                    {
                        max_compute_perf  = compute_perf;
                        max_perf_device   = current_device;
                    }
                }
                else
                {
                    max_compute_perf  = compute_perf;
                    max_perf_device   = current_device;
                }
            }
        }
        ++current_device;
    }
    return max_perf_device;
}

void findCudaDevice()
{
    cudaDeviceProp deviceProp;
    int devID = 0;
    devID = getBestDevice();
    CUDA_CHECK_RETURN(cudaSetDevice(devID));
    CUDA_CHECK_RETURN(cudaGetDeviceProperties(&deviceProp, devID));
    printf("> Using CUDA device [%d]: %s\n", devID, deviceProp.name);
}

/*
 * Function that starts GPU Kernel.
 */
void launchKernel(vector<char> &inputVector, int blocksPerGrid, int threadsPerBlock, unsigned int allHits, int topHitsNum, char **results){
	findCudaDevice(); //choosing the fastest (GFLOPS) card on the system
	void *vecKernel = NULL;
	void *resultsKernel = NULL;

	CUDA_CHECK_RETURN(cudaDeviceReset());
	//cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); //TODO: we can increase L1 cache if we don't use shared mem
	int resultsInBytes = (allHits)*(sizeof(track) + 2*sizeof(int));   //Allocating memory for as many tracks as hits...we can allocate 1/3 of it but then we must know all the offeset between events
	CUDA_CHECK_RETURN(cudaMalloc((void**) &vecKernel, inputVector.size()));
	CUDA_CHECK_RETURN(cudaMemcpy(vecKernel, &inputVector[0], inputVector.size(), cudaMemcpyHostToDevice));
	CUDA_CHECK_RETURN(cudaMalloc((void**) &resultsKernel, resultsInBytes));

	searchByPair<<<blocksPerGrid, threadsPerBlock>>>(vecKernel,resultsKernel,topHitsNum);
//	searchByPair<<<1,4>>>(vecKernel,resultsKernel, allHits); //testing

	CUDA_CHECK_RETURN(cudaThreadSynchronize());	// Wait for the GPU launched work to complete
	CUDA_CHECK_RETURN(cudaGetLastError());

	*results = (char*)malloc(resultsInBytes);  //TODO: release 'results' memory
	CUDA_CHECK_RETURN(cudaMemcpy(*results, resultsKernel, resultsInBytes, cudaMemcpyDeviceToHost));
	CUDA_CHECK_RETURN(cudaFree((void*) vecKernel));
	CUDA_CHECK_RETURN(cudaFree((void*) resultsKernel));
	CUDA_CHECK_RETURN(cudaDeviceReset());

}

	.file	"tmpxft_001fe6d1_00000000-6_SearchKernel.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4150:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4150:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z5zBeamP5track
	.type	_Z5zBeamP5track, @function
_Z5zBeamP5track:
.LFB4129:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4129:
	.size	_Z5zBeamP5track, .-_Z5zBeamP5track
	.globl	_Z5r2AtZdP5track
	.type	_Z5r2AtZdP5track, @function
_Z5r2AtZdP5track:
.LFB4130:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4130:
	.size	_Z5r2AtZdP5track, .-_Z5r2AtZdP5track
	.globl	_Z5solveP5track
	.type	_Z5solveP5track, @function
_Z5solveP5track:
.LFB4131:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4131:
	.size	_Z5solveP5track, .-_Z5solveP5track
	.globl	_Z6addHitP5trackiPiPdS2_S2_S2_
	.type	_Z6addHitP5trackiPiPdS2_S2_S2_, @function
_Z6addHitP5trackiPiPdS2_S2_S2_:
.LFB4132:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4132:
	.size	_Z6addHitP5trackiPiPdS2_S2_S2_, .-_Z6addHitP5trackiPiPdS2_S2_S2_
	.globl	_Z8setTrackP5trackiiPiPdS2_S2_S2_
	.type	_Z8setTrackP5trackiiPiPdS2_S2_S2_, @function
_Z8setTrackP5trackiiPiPdS2_S2_S2_:
.LFB4133:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4133:
	.size	_Z8setTrackP5trackiiPiPdS2_S2_S2_, .-_Z8setTrackP5trackiiPiPdS2_S2_S2_
	.section	.rodata
	.align 32
.LC0:
	.long	16
	.long	8
	.long	17
	.long	8
	.long	18
	.long	8
	.long	19
	.long	8
	.long	32
	.long	32
	.long	33
	.long	48
	.long	48
	.long	192
	.long	53
	.long	192
	.section	.text._Z13getBestDevicev,"axG",@progbits,_Z13getBestDevicev,comdat
	.weak	_Z13getBestDevicev
	.type	_Z13getBestDevicev, @function
_Z13getBestDevicev:
.LFB4145:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	xorl	%ebx, %ebx
	subq	$1120, %rsp
	.cfi_def_cfa_offset 1168
	movq	%fs:40, %rax
	movq	%rax, 1112(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rdi
	xorl	%eax, %eax
	leaq	80(%rsp), %r12
	movl	%eax, 12(%rsp)
	call	cudaGetDeviceCount@PLT
.L13:
	cmpl	%ebx, 12(%rsp)
	jle	.L31
	movl	%ebx, %esi
	movq	%r12, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	cmpl	$2, 484(%rsp)
	je	.L14
	movl	440(%rsp), %eax
	leal	-1(%rax), %edx
	cmpl	$9997, %edx
	ja	.L14
	cmpl	%eax, %ebp
	cmovl	%eax, %ebp
.L14:
	incl	%ebx
	jmp	.L13
.L31:
	xorl	%r12d, %r12d
	xorl	%r13d, %r13d
	xorl	%ebx, %ebx
	movabsq	$42945378002703, %r14
.L16:
	cmpl	%ebx, 12(%rsp)
	jle	.L32
	leaq	80(%rsp), %rdi
	movl	%ebx, %esi
	call	cudaGetDeviceProperties_v2@PLT
	cmpl	$2, 484(%rsp)
	je	.L17
	movl	$1, %eax
	cmpq	%r14, 440(%rsp)
	je	.L18
	movl	440(%rsp), %edx
	leaq	16(%rsp), %rdi
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	movl	$16, %ecx
	rep movsl
	sall	$4, %edx
	leaq	16(%rsp), %rsi
	addl	444(%rsp), %edx
.L19:
	movl	(%rsi,%rax,8), %ecx
	cmpl	$-1, %ecx
	je	.L33
	leaq	1(%rax), %rdi
	cmpl	%ecx, %edx
	jne	.L27
	cltq
	movl	20(%rsp,%rax,8), %eax
	jmp	.L18
.L27:
	movq	%rdi, %rax
	jmp	.L19
.L33:
	movl	$192, %eax
.L18:
	imull	468(%rsp), %eax
	imull	428(%rsp), %eax
	cmpl	%eax, %r13d
	jl	.L22
.L17:
	movl	%r13d, %eax
	jmp	.L23
.L22:
	cmpl	$2, %ebp
	jle	.L28
	cmpl	%ebp, 440(%rsp)
	jne	.L17
.L28:
	movl	%ebx, %r12d
.L23:
	incl	%ebx
	movl	%eax, %r13d
	jmp	.L16
.L32:
	movq	1112(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$1120, %rsp
	.cfi_def_cfa_offset 48
	movl	%r12d, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4145:
	.size	_Z13getBestDevicev, .-_Z13getBestDevicev
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/dcampora/gpu-manager/master/gpu-ip-example/GPUSearch/src/SearchKernel.cu"
.LC2:
	.string	"Error %s at line %d in file %s\n"
.LC3:
	.string	"> Using CUDA device [%d]: %s\n"
	.text
	.globl	_Z14findCudaDevicev
	.type	_Z14findCudaDevicev, @function
_Z14findCudaDevicev:
.LFB4146:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$1048, %rsp
	.cfi_def_cfa_offset 1072
	movq	%fs:40, %rax
	movq	%rax, 1032(%rsp)
	xorl	%eax, %eax
	call	_Z13getBestDevicev
	movl	%eax, %edi
	movl	%eax, %ebx
	call	cudaSetDevice@PLT
	testl	%eax, %eax
	je	.L35
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$588, %r8d
	movq	%rax, %rcx
	jmp	.L39
.L35:
	movq	%rsp, %rbp
	movl	%ebx, %esi
	movq	%rbp, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L36
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$589, %r8d
	movq	%rax, %rcx
.L39:
	movq	stderr(%rip), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L36:
	xorl	%eax, %eax
	movq	%rbp, %rcx
	movl	%ebx, %edx
	movl	$2, %edi
	leaq	.LC3(%rip), %rsi
	call	__printf_chk@PLT
	movq	1032(%rsp), %rax
	subq	%fs:40, %rax
	je	.L37
	call	__stack_chk_fail@PLT
.L37:
	addq	$1048, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4146:
	.size	_Z14findCudaDevicev, .-_Z14findCudaDevicev
	.globl	_Z35__device_stub__Z12searchByPairPvS_iPvS_i
	.type	_Z35__device_stub__Z12searchByPairPvS_iPvS_i, @function
_Z35__device_stub__Z12searchByPairPvS_iPvS_i:
.LFB4172:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z12searchByPairPvS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L40:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L42
	call	__stack_chk_fail@PLT
.L42:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4172:
	.size	_Z35__device_stub__Z12searchByPairPvS_iPvS_i, .-_Z35__device_stub__Z12searchByPairPvS_iPvS_i
	.globl	_Z12searchByPairPvS_i
	.type	_Z12searchByPairPvS_i, @function
_Z12searchByPairPvS_i:
.LFB4173:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z12searchByPairPvS_iPvS_i
	.cfi_endproc
.LFE4173:
	.size	_Z12searchByPairPvS_i, .-_Z12searchByPairPvS_i
	.globl	_Z12launchKernelRSt6vectorIcSaIcEEiijiPPc
	.type	_Z12launchKernelRSt6vectorIcSaIcEEiijiPPc, @function
_Z12launchKernelRSt6vectorIcSaIcEEiijiPPc:
.LFB4147:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%edx, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%esi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%r8d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%r9, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Z14findCudaDevicev
	xorl	%eax, %eax
	movq	%rax, (%rsp)
	movq	%rax, 8(%rsp)
	call	cudaDeviceReset@PLT
	testl	%eax, %eax
	je	.L46
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$601, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L46:
	movq	8(%rbp), %rsi
	movq	%rsp, %rdi
	subq	0(%rbp), %rsi
	call	cudaMalloc@PLT
	imull	$136, %ebx, %ebx
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L47
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$604, %r8d
	movq	%rax, %rcx
.L59:
	movq	stderr(%rip), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L47:
	movq	0(%rbp), %rsi
	movq	8(%rbp), %rdx
	movl	$1, %ecx
	movq	(%rsp), %rdi
	subq	%rsi, %rdx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L48
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$605, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L48:
	movslq	%ebx, %rbx
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L49
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$606, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L49:
	movabsq	$4294967297, %rax
	movl	%r15d, 28(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rax, 32(%rsp)
	movl	36(%rsp), %ecx
	movq	%rax, 20(%rsp)
	movq	28(%rsp), %rdx
	movl	%r14d, 16(%rsp)
	movl	24(%rsp), %esi
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L50
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	%r13d, %edx
	call	_Z35__device_stub__Z12searchByPairPvS_iPvS_i
.L50:
	call	cudaThreadSynchronize@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L51
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$611, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L51:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L52
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$612, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L52:
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, (%r12)
	movq	%rax, %rdi
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L53
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$615, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L53:
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L54
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$616, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L54:
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L55
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$617, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L55:
	call	cudaDeviceReset@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L45
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %r9
	movl	$618, %r8d
	movq	%rax, %rcx
	jmp	.L59
.L45:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L57
	call	__stack_chk_fail@PLT
.L57:
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4147:
	.size	_Z12launchKernelRSt6vectorIcSaIcEEiijiPPc, .-_Z12launchKernelRSt6vectorIcSaIcEEiijiPPc
	.section	.rodata.str1.1
.LC4:
	.string	"_Z12searchByPairPvS_i"
.LC5:
	.string	"m_maxXSlope"
.LC6:
	.string	"m_maxYSlope"
.LC7:
	.string	"m_maxZForRBeamCut"
.LC8:
	.string	"m_maxR2Beam"
.LC9:
	.string	"m_maxMissed"
.LC10:
	.string	"m_extraTol"
.LC11:
	.string	"m_maxChi2ToAdd"
.LC12:
	.string	"m_maxChi2SameSensor"
.LC13:
	.string	"m_maxChi2Short"
.LC14:
	.string	"m_maxChi2PerHit"
.LC15:
	.string	"m_sensNum"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4175:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC4(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z12searchByPairPvS_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movl	$8, %r9d
	leaq	_ZL11m_maxXSlope(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL11m_maxYSlope(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL17m_maxZForRBeamCut(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL11m_maxR2Beam(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC9(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL11m_maxMissed(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC10(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL10m_extraTol(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL14m_maxChi2ToAdd(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC12(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL19m_maxChi2SameSensor(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC13(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL14m_maxChi2Short(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL15m_maxChi2PerHit(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC15(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL9m_sensNum(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4175:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.local	_ZL9m_sensNum
	.comm	_ZL9m_sensNum,4,4
	.local	_ZL15m_maxChi2PerHit
	.comm	_ZL15m_maxChi2PerHit,8,8
	.local	_ZL14m_maxChi2Short
	.comm	_ZL14m_maxChi2Short,8,8
	.local	_ZL19m_maxChi2SameSensor
	.comm	_ZL19m_maxChi2SameSensor,8,8
	.local	_ZL14m_maxChi2ToAdd
	.comm	_ZL14m_maxChi2ToAdd,8,8
	.local	_ZL10m_extraTol
	.comm	_ZL10m_extraTol,8,8
	.local	_ZL11m_maxMissed
	.comm	_ZL11m_maxMissed,4,4
	.local	_ZL11m_maxR2Beam
	.comm	_ZL11m_maxR2Beam,8,8
	.local	_ZL17m_maxZForRBeamCut
	.comm	_ZL17m_maxZForRBeamCut,8,8
	.local	_ZL11m_maxYSlope
	.comm	_ZL11m_maxYSlope,8,8
	.local	_ZL11m_maxXSlope
	.comm	_ZL11m_maxXSlope,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

//pragma once
#ifndef SEARCHKERNEL_H_
#define SEARCHKERNEL_H_
#include <vector>

void launchKernel(std::vector<char> &inputVector, int eventsNumber, int *events, unsigned int allHits, char **results);
inline int getBestDevice();
void findCudaDevice();

#endif /* SEARCHKERNEL_H_ */

	.text
	.file	"SearchKernel.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
