
// 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: mutation <<<N_BLOCK,BLOCK_LENGTH>>> (pgpu,randomPM,randomPoint,REAL_LEN,MASK,PROB_MUT);
// N_BLOCK: number of blocks, such that N_BLOCK * BLOCK_LENGTH = POP_SIZE.
// BLOCK_LENGTH: length of each block, such that N_BLOCK * BLOCK_LENGTH = POP_SIZE.
// pgpu: pointer to global memory that stores the population. 
// randomPM: pointer to global memory that stores the random values for mutation.
// randomPoint: pointer to global memory that stores the random points for mutation.
// REAL_LEN: real length of the chromosome of datatype Data. 
// REAL_LEN = (CHROM_LEN%DataSize)==0?CHROM_LEN/DataSize:(CHROM_LEN/DataSize + 1) 
// MASK: A bit mask of the first bit. MASK = (Data)pow(2.0, (int)DataSize-1)
// PROB_MUT: mutation probability. 

__global__ void mutation(Data *pop, float *randomPM, int *randomPoint, int length, Data mask, float PROB_MUT){
	
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float pm = randomPM[idx];   // value for mutation
    int pnt = randomPoint[idx]; // mutation point
	
    if (pm <= PROB_MUT) {
          int word = pnt/DataSize;
	  int wordPos = pnt % DataSize;
          Data aux1 = pop[idx*length+word];
	  Data aux2 = mask >> wordPos;
	  pop[idx*length+word] = aux1 ^ aux2;
    }
 
}


