
// 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: fitness <<<POP_SIZE,MAX_THREADS_PER_BLOCK>>> (pgpu,fitgpu,REAL_LEN,MASK,CHROM_LEN);
// POP_SIZE: number of individuals of the population.
// MAX_THREADS_PER_BLOCK: constant with the maximum number of threads per block. 
// CHROM_LEN: length of the chromosome.
// pgpu: pointer to the population. 
// fitgpu: pointer to an auxiliary structure to store the fitness.
// MASK: A bit mask of the first bit. MASK = (Data)pow(2.0, (int)DataSize-1)
// REAL_LEN: real length of the chromosome of datatype Data. 
// REAL_LEN = (CHROM_LEN%DataSize)==0?CHROM_LEN/DataSize:(CHROM_LEN/DataSize + 1) 

__global__ void fitness(Data * pop, int * fit, int realLength, Data mask, int length) {
    
    __shared__ int partial[MAX_THREADS_PER_BLOCK];  // array to store partial fitness
    int idx = blockIdx.x;  // the number of the individual is indicated by the block number
    partial[threadIdx.x] = 0; // each array position is initialized in 0

    int i,j;
    // each thread adds in partial its corresponding values
    Data aux, shift, value;
    for(i=threadIdx.x;i<realLength;i=i+MAX_THREADS_PER_BLOCK) {
         aux = pop[idx*realLength+i];
	 shift = mask;
	 for(j=0;j<DataSize && (i*DataSize+j<length);j=j+1) {
	       value = (aux & shift)==0?0:1;
	       shift = shift >> 1;
	       partial[threadIdx.x] = partial[threadIdx.x] + value;
	 }
    }

    // reduction algorithm to add the partial fitness values
    __syncthreads();
    i = MAX_THREADS_PER_BLOCK/2;
    while (i!=0) {
       if (threadIdx.x < i)
	  partial[threadIdx.x] = partial[threadIdx.x] + partial[threadIdx.x + i]; 
       __syncthreads();
       i = i/2;
     }
    
    // finally thread 0 writes the fitness value in global memory
    if (threadIdx.x == 0)
	fit[idx] = partial[0];	
}