
// The datatype Data should be defined as unsigned char, unsigned short, unsigned int or unsigned long.
typedef ... Data;
// The constant DataSize should be defined as 8 (if Data is unsigned char), 16 (if Data is unsigned short), 
// 32 (if Data is unsigned int) or 64 (if Data is unsigned long).
#define DataSize ...

// Kernel invocation: dpx <<POP_SIZE/2,MAX_THREADS_PER_BLOCK>>> (oldpop,newpop,positions,randomPC,randomPoint,
// REAL_LEN,PROB_CROSS); 
// POP_SIZE: number of individuals of the population.
// MAX_THREADS_PER_BLOCK: constant with the maximum number of threads per block. 
// oldpop: pointer to global memory that stores the old population. 
// newpop: pointer to global memory that stores the new population.
// positions: pointer to global memory that stores the indexes of individuals for crossover
// randomPC: pointer to global memory that stores the random values for crossover.
// randomPoint: pointer to global memory that stores the random points for crossover.
// REAL_LEN: real length of the chromosome of datatype Data. 
// REAL_LEN = (CHROM_LEN%DataSize)==0?CHROM_LEN/DataSize:(CHROM_LEN/DataSize + 1) 
// PROB_CROSS: crossover probability. 

__global__ void dpx(Data *pop, Data *npop, int *pos, float *randomPC, int *randomPoint, int length, float PROB_CROSS){

   int idx = blockIdx.x;  // the number of the individual is indicated by the block number
   int ind1 = pos[2*idx];  // index of first individual for crossover
   int ind2 = pos[2*idx+1]; // index of second individual for crossover
   float pc = randomPC[idx]; // value for crossover
   int pnt1 = randomPoint[2*idx]; // first crossover point
   int pnt2 = randomPoint[2*idx+1]; // second crossover point. pnt2 != pnt1
   int word1 = pnt1/DataSize;     // word of the first crossover point
   unsigned int wordPos1 = pnt1 % DataSize; // position in the word of the first crossover point 
   int word2 = pnt2/DataSize;     // word of the second crossover point
   unsigned int wordPos2 = pnt2 % DataSize; // position in the word of the second crossover point 

   int i;

   if (pc <= PROB_CROSS) {
       // Cross individuals
       if (word1!=word2) { 
          unsigned int move;
          for(i=threadIdx.x;i<length;i=i+MAX_THREADS_PER_BLOCK) {
             if (i<word1 || i>word2 || (i==word2 && wordPos2==0)) {
	        //copy word from parent ind1 to child 2*idx
	        npop[2*idx*length+i] = pop[ind1*length+i];
		//copy word from parent ind2 to child 2*idx + 1
 		npop[(2*idx+1)*length+i] = pop[ind2*length+i];
	     } else if (i==word1 && wordPos1 > 0) {
                 // the word has to be shifted  
       		 move = DataSize -  wordPos1;
    		 npop[2*idx*length+i] = ((Data)(pop[ind1*length+i] >> move) << move) | ((Data)(pop[ind2*length+i] << wordPos1) >>  wordPos1);
		 npop[(2*idx+1)*length+i] = ((Data)(pop[ind2*length+i] >> move) << move) | ((Data)(pop[ind1*length+i] << wordPos1) >> wordPos1);
	     } else if (i==word2 && wordPos2 > 0) {
		 // the word has to be shifted
                 move = DataSize -  wordPos2;
		 npop[2*idx*length+i] = ((Data)(pop[ind2*length+i] >> move) << move) | ((Data)(pop[ind1*length+i] << wordPos2) >>  wordPos2);
		 npop[(2*idx+1)*length+i] = ((Data)(pop[ind1*length+i] >> move) << move) | ((Data)(pop[ind2*length+i] << wordPos2) >>  wordPos2);
	     } else {
                 //copy word from parent ind2 to child 2*idx	       		    
                 npop[2*idx*length+i] = pop[ind2*length+i];
		 //copy word from parent ind1 to child 2*idx + 1
	         npop[(2*idx+1)*length+i] = pop[ind1*length+i];
	     }
          }
       } else { // wordPos1 != wordPos2 since pnt2 != pnt1
              for(i=threadIdx.x;i<length;i=i+MAX_THREADS_PER_BLOCK) {
                  if (i<word1) {
                    //copy word from parent ind1 to child 2*idx
	            npop[2*idx*length+i] = pop[ind1*length+i];
		    //copy word from parent ind2 to child 2*idx + 1
 		    npop[(2*idx+1)*length+i] = pop[ind2*length+i];
	          } else if (i==word1) {
                    // the word has to be shifted 
                    unsigned int move = DataSize - wordPos1; 
	            unsigned int move2 = DataSize - wordPos2;
	       	    Data rest1 = ((Data)(pop[ind2*length+i] << wordPos1) >>  wordPos1);
	       	    Data rest2 = ((Data)(pop[ind1*length+i] << wordPos1) >>  wordPos1);
                    if (wordPos1!=0) {
                        npop[2*idx*length+i] = ((Data)(pop[ind1*length+i] >> move) << move) | ((Data)(rest1 >> move2) << move2) | ((Data)(rest2 << wordPos2) >>  wordPos2);
		        npop[(2*idx+1)*length+i] = ((Data)(pop[ind2*length+i] >> move) << move) | ((Data)(rest2 >> move2) << move2) | ((Data)(rest1 << wordPos2) >>  wordPos2);
		     } else {
 		        npop[2*idx*length+i] = ((Data)(rest1 >> move2) << move2) | ((Data)(rest2 << wordPos2) >>  wordPos2);
  			npop[(2*idx+1)*length+i] = ((Data)(rest2 >> move2) << move2) | ((Data)(rest1 << wordPos2) >>  wordPos2);
                     }
	          } else {
                      //copy word from parent ind2 to child 2*idx	       		    
                      npop[2*idx*length+i] = pop[ind2*length+i];
		      //copy word from parent ind1 to child 2*idx + 1
	              npop[(2*idx+1)*length+i] = pop[ind1*length+i];
	          }
              }
      }
   } else {
       // Copy individuals
       for(i=threadIdx.x;i<length;i=i+MAX_THREADS_PER_BLOCK) {
         //copy word from parent ind1 to child 2*idx
	 npop[(2*idx)*length+i] = pop[ind1*length+i];
	 //copy word from parent ind2 to child 2*idx + 1
	 npop[(2*idx+1)*length+i] = pop[ind2*length+i];
     }
   }
}
