
// 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: spx <<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 spx(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 pnt = randomPoint[idx]; // crossover point
  int i;
	
  if (pc <= PROB_CROSS) {
     // Cross individuals
     for(i=threadIdx.x;i<length;i=i+MAX_THREADS_PER_BLOCK) {
	  if (i<(pnt/DataSize)) {	  
              //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==(pnt/DataSize) && pnt % DataSize > 0) {
	         unsigned int word = pnt/DataSize;
	         unsigned int move = pnt % DataSize;
	         unsigned int moveComp = DataSize - move;
                 npop[2*idx*length+word] = ((Data)(pop[ind1*length+word] >> moveComp) << moveComp) | ((Data)(pop[ind2*length+word] << move) >>  move);
                 npop[(2*idx+1)*length+word] = ((Data)(pop[ind2*length+word] >> moveComp) << moveComp) | ((Data)(pop[ind1*length+word] << move) >>  move);	  
	  } 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];
      }
 }

}
