git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2083 331e1502-861f-0410-8da2-ba01fb791d7f

This commit is contained in:
boufaras 2011-01-28 09:51:56 +00:00
commit 0ab74e3bb2

View file

@ -38,7 +38,7 @@
#include <eval/moCudakernelEval.h> #include <eval/moCudakernelEval.h>
/** /**
* class for the cuda evaluation * class for the K-swap neighborhood evaluation
*/ */
template<class Neighbor, class IncrementEval> template<class Neighbor, class IncrementEval>
@ -75,7 +75,6 @@ public:
moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) { moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
mutex = false; mutex = false;
mutex_kswap=false; mutex_kswap=false;
compt=0;
} }
/** /**
* Destructor * Destructor
@ -99,7 +98,7 @@ public:
/** /**
* Compute fitness for all solution neighbors in device * Compute fitness for all solution neighbors in device
* @param _sol the solution which generate the neighborhood * @param _sol the solution which generate the neighborhood
* @param _mapping the array of indices mapping * @param _mapping the array of mapping indexes for K-swap neighborhood
* @param _Kswap the number of swap * @param _Kswap the number of swap
*/ */
@ -110,6 +109,7 @@ public:
// Get Current solution fitness // Get Current solution fitness
Fitness fitness = _sol.fitness(); Fitness fitness = _sol.fitness();
//Case of Permutation //Case of Permutation
if (_Kswap == 1) { if (_Kswap == 1) {
if (!mutex) { if (!mutex) {
@ -117,6 +117,7 @@ public:
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
mutex = true; mutex = true;
} }
//Copy the solution vector from the host to device //Copy the solution vector from the host to device
cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T), cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
cudaMemcpyHostToDevice); cudaMemcpyHostToDevice);
@ -130,6 +131,7 @@ public:
//Case Kswap //Case Kswap
else if (_Kswap > 1) { else if (_Kswap > 1) {
if (!mutex_kswap) { if (!mutex_kswap) {
vect = new T[neighborhoodSize * _size]; vect = new T[neighborhoodSize * _size];
//Allocate the space for set of solution in the device global memory //Allocate the space for set of solution in the device global memory
cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize
@ -138,6 +140,7 @@ public:
cudaMalloc((void**) &device_tmp.vect, neighborhoodSize cudaMalloc((void**) &device_tmp.vect, neighborhoodSize
* sizeof(T)); * sizeof(T));
mutex_kswap = true; mutex_kswap = true;
} }
for (int i = 0; i < neighborhoodSize; i++) { for (int i = 0; i < neighborhoodSize; i++) {
@ -145,19 +148,17 @@ public:
vect[j + i * _size] = _sol.vect[j]; vect[j + i * _size] = _sol.vect[j];
} }
} }
//Copy the set of solution from the host to device //Copy the set of solution from the host to device
cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size
* sizeof(T), cudaMemcpyHostToDevice); * sizeof(T), cudaMemcpyHostToDevice);
//Launch the Kernel to compute all Kswap neighbors fitness //Launch the Kernel to compute all Kswap neighbors fitness
kernelKswap<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_setSolution,device_tmp,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap,_size); kernelKswap<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_setSolution,device_tmp,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap,_size);
//Copy the result from device to host //Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost); * sizeof(Fitness), cudaMemcpyDeviceToHost);
std::cout<<" Iteration number "<<compt<<" "<<std::endl;
/* for(unsigned int i=0;i<neighborhoodSize;i++)
std::cout<<" "<<host_FitnessArray[i]<<" ";//<<std::endl;*/
compt++;
} }
} }
@ -165,7 +166,7 @@ public:
/** /**
* Compute fitness for all solution neighbors(K-flip of binary solution) in device * Compute fitness for all solution neighbors(K-flip of binary solution) in device
* @param _sol the solution which generate the neighborhood * @param _sol the solution which generate the neighborhood
* @param _mapping the array of indices mapping * @param _mapping the array of mapping indexes for k-flip neighborhood
* @param _Kflip the number of flip to do * @param _Kflip the number of flip to do
*/ */
@ -176,17 +177,20 @@ public:
// Get Current solution fitness // Get Current solution fitness
Fitness fitness = _sol.fitness(); Fitness fitness = _sol.fitness();
if (!mutex) { if (!mutex) {
//Allocate the space for solution in the device global memory //Allocate the space for solution in the device global memory
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
mutex = true; mutex = true;
} }
//Copy the solution vector from the host to device //Copy the solution vector from the host to device
cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T), cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
cudaMemcpyHostToDevice); cudaMemcpyHostToDevice);
//Launch the Kernel to compute all flip neighbors fitness //Launch the Kernel to compute all flip neighbors fitness
kernelKflip<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kflip); kernelKflip<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kflip);
//Copy the result from device to host //Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost); * sizeof(Fitness), cudaMemcpyDeviceToHost);
@ -203,7 +207,6 @@ protected:
T * vect; T * vect;
bool mutex_kswap; bool mutex_kswap;
bool mutex; bool mutex;
unsigned int compt;
}; };
#endif #endif