
// Kernel invocation: fitness <<<POP_SIZE,MAX_THREADS_PER_BLOCK>>> (pgpu,fitgpu,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 global memory that stores the population. 
// fitgpu: pointer to global memory that stores the fitness values.

__global__ void fitness(bool * pop, int * fit, 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;
    // each thread adds in partial its corresponding values
    for(i=threadIdx.x;i<length;i=i+MAX_THREADS_PER_BLOCK) {
        partial[threadIdx.x] = partial[threadIdx.x] + pop[idx*length+i];
    }
	
    // 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];	
}
