nojhan/paradiseo
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity

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

Browse source
This commit is contained in:
boufaras 2011-02-10 13:43:14 +00:00
commit 87dd3eaa92
1 changed files with 190 additions and 191 deletions
Download patch file Download diff file Expand all files Collapse all files

337
ParadisEO-GPU/src/eval/moCudaKswapEval.h
View file

@ -1,36 +1,36 @@
/* /*
<moCudaKswapEval.h> <moCudaKswapEval.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Karima Boufaras, Thé Van LUONG Karima Boufaras, Thé Van LUONG
This software is governed by the CeCILL license under French law and This software is governed by the CeCILL license under French law and
abiding by the rules of distribution of free software. You can use, abiding by the rules of distribution of free software. You can use,
modify and/ or redistribute the software under the terms of the CeCILL modify and/ or redistribute the software under the terms of the CeCILL
license as circulated by CEA, CNRS and INRIA at the following URL license as circulated by CEA, CNRS and INRIA at the following URL
"http://www.cecill.info". "http://www.cecill.info".
As a counterpart to the access to the source code and rights to copy, As a counterpart to the access to the source code and rights to copy,
modify and redistribute granted by the license, users are provided only modify and redistribute granted by the license, users are provided only
with a limited warranty and the software's author, the holder of the with a limited warranty and the software's author, the holder of the
economic rights, and the successive licensors have only limited liability. economic rights, and the successive licensors have only limited liability.
In this respect, the user's attention is drawn to the risks associated In this respect, the user's attention is drawn to the risks associated
with loading, using, modifying and/or developing or reproducing the with loading, using, modifying and/or developing or reproducing the
software by the user in light of its specific status of free software, software by the user in light of its specific status of free software,
that may mean that it is complicated to manipulate, and that also that may mean that it is complicated to manipulate, and that also
therefore means that it is reserved for developers and experienced therefore means that it is reserved for developers and experienced
professionals having in-depth computer knowledge. Users are therefore professionals having in-depth computer knowledge. Users are therefore
encouraged to load and test the software's suitability as regards their encouraged to load and test the software's suitability as regards their
requirements in conditions enabling the security of their systems and/or requirements in conditions enabling the security of their systems and/or
data to be ensured and, more generally, to use and operate it in the data to be ensured and, more generally, to use and operate it in the
same conditions as regards security. same conditions as regards security.
The fact that you are presently reading this means that you have had The fact that you are presently reading this means that you have had
knowledge of the CeCILL license and that you accept its terms. knowledge of the CeCILL license and that you accept its terms.
ParadisEO WebSite : http://paradiseo.gforge.inria.fr ParadisEO WebSite : http://paradiseo.gforge.inria.fr
Contact: paradiseo-help@lists.gforge.inria.fr Contact: paradiseo-help@lists.gforge.inria.fr
*/ */
#ifndef moCudaKswapEval_H #ifndef moCudaKswapEval_H
#define moCudaKswapEval_H #define moCudaKswapEval_H
@ -42,172 +42,171 @@
*/ */
template<class Neighbor, class IncrementEval> template<class Neighbor, class IncrementEval>
class moCudaKswapEval: public moCudaEval<Neighbor> { class moCudaKswapEval: public moCudaEval<Neighbor> {
public: public:
/** /**
* Define type of a solution corresponding to Neighbor * Define type of a solution corresponding to Neighbor
*/ */
typedef typename Neighbor::EOT EOT; typedef typename Neighbor::EOT EOT;
/** /**
* Define type of a vector corresponding to Solution * Define type of a vector corresponding to Solution
*/ */
typedef typename EOT::ElemType T; typedef typename EOT::ElemType T;
/** /**
* Define type of a fitness corresponding to Solution * Define type of a fitness corresponding to Solution
*/ */
typedef typename EOT::Fitness Fitness; typedef typename EOT::Fitness Fitness;
using moCudaEval<Neighbor>::neighborhoodSize; using moCudaEval<Neighbor>::neighborhoodSize;
using moCudaEval<Neighbor>::host_FitnessArray; using moCudaEval<Neighbor>::host_FitnessArray;
using moCudaEval<Neighbor>::device_FitnessArray; using moCudaEval<Neighbor>::device_FitnessArray;
using moCudaEval<Neighbor>::device_solution; using moCudaEval<Neighbor>::device_solution;
using moCudaEval<Neighbor>::kernel_Dim; using moCudaEval<Neighbor>::kernel_Dim;
/** /**
* Constructor * Constructor
* @param _neighborhoodSize the size of the neighborhood * @param _neighborhoodSize the size of the neighborhood
* @param _incrEval the incremental evaluation * @param _incrEval the incremental evaluation
*/ */
moCudaKswapEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) : moCudaKswapEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) :
moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) { moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
mutex = false; mutex = false;
mutex_kswap = false; mutex_kswap = false;
}
/**
* Destructor
*/
~moCudaKswapEval() {
if (mutex_kswap) {
cudaFree(&device_setSolution);
cudaFree(&device_tmp);
delete[] vect;
}
}
/**
* Compute fitness for all solution neighbors in device without specific mapping
* @param _sol the solution which generate the neighborhood
*/
virtual void neighborhoodEval(EOT & _sol) {
}
/**
* Compute fitness for all solution neighbors in device with K-swap mapping
* @param _sol the solution which generate the neighborhood
* @param _mapping the array of mapping indexes for K-swap neighborhood
* @param _Kswap the number of swap
*/
void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping, unsigned _Kswap) {
// the solution size
unsigned _size = _sol.size();
// Get Current solution fitness
Fitness fitness = _sol.fitness();
//Case of Permutation
if (_Kswap == 1) {
if (!mutex) {
//Allocate the space for solution in the device global memory
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
mutex = true;
} }
/** //Copy the solution vector from the host to device
* Destructor cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
*/ cudaMemcpyHostToDevice);
~moCudaKswapEval() {
if (mutex_kswap) { //Launch the Kernel to compute all permutation neighbors fitness
cudaFree(&device_setSolution); kernelPermutation<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_size);
cudaFree(&device_tmp); //Copy the result from device to host
delete[] vect; cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
} * sizeof(Fitness), cudaMemcpyDeviceToHost);
}
//Case Kswap
else if (_Kswap > 1) {
if (!mutex_kswap) {
vect = new T[neighborhoodSize * _size];
//Allocate the space for set of solution in the device global memory
cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize
* _size * sizeof(T));
//Allocate the space to save temporary EOT element to swap
cudaMalloc((void**) &device_tmp.vect, neighborhoodSize
* sizeof(T));
mutex_kswap = true;
} }
/** for (int i = 0; i < neighborhoodSize; i++) {
* Compute fitness for all solution neighbors in device without specific mapping for (int j = 0; j < _size; j++) {
* @param _sol the solution which generate the neighborhood vect[j + i * _size] = _sol.vect[j];
*/ }
virtual void neighborhoodEval(EOT & _sol) {
} }
/** //Copy the set of solution from the host to device
* Compute fitness for all solution neighbors in device with K-swap mapping cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size
* @param _sol the solution which generate the neighborhood * sizeof(T), cudaMemcpyHostToDevice);
* @param _mapping the array of mapping indexes for K-swap neighborhood
* @param _Kswap the number of swap
*/
void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping, unsigned _Kswap) { //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);
// the solution size //Copy the result from device to host
unsigned _size = _sol.size(); cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
// Get Current solution fitness }
Fitness fitness = _sol.fitness(); }
//Case of Permutation /**
if (_Kswap == 1) { * Compute fitness for all solution neighbors(K-flip of binary solution) in device
if (!mutex) { * @param _sol the solution which generate the neighborhood
//Allocate the space for solution in the device global memory * @param _mapping the array of mapping indexes for k-flip neighborhood
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); * @param _Kflip the number of flip to do
mutex = true; */
}
//Copy the solution vector from the host to device void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping,
cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T), unsigned _Kflip) {
cudaMemcpyHostToDevice);
//Launch the Kernel to compute all permutation neighbors fitness // the solution size
kernelPermutation<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_size); unsigned _size = _sol.size();
//Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
}
//Case Kswap
else if (_Kswap > 1) {
if (!mutex_kswap) {
vect = new T[neighborhoodSize * _size];
//Allocate the space for set of solution in the device global memory
cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize
* _size * sizeof(T));
//Allocate the space to save temporary EOT element to swap
cudaMalloc((void**) &device_tmp.vect, neighborhoodSize
* sizeof(T));
mutex_kswap = true;
}
for (int i = 0; i < neighborhoodSize; i++) { // Get Current solution fitness
for (int j = 0; j < _size; j++) { Fitness fitness = _sol.fitness();
vect[j + i * _size] = _sol.vect[j]; if (!mutex) {
} //Allocate the space for solution in the device global memory
} cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
mutex = true;
}
//Copy the set of solution from the host to device //Copy the solution vector from the host to device
cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
* sizeof(T), cudaMemcpyHostToDevice); cudaMemcpyHostToDevice);
//Launch the Kernel to compute all Kswap neighbors fitness //Launch the Kernel to compute all flip neighbors fitness
kernelKswap<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_setSolution,device_tmp,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap,_size); 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);
}
} protected:
}
/** IncrementEval & incrEval;
* Compute fitness for all solution neighbors(K-flip of binary solution) in device //NeighborhoodSize copy of solution
* @param _sol the solution which generate the neighborhood EOT device_setSolution;
* @param _mapping the array of mapping indexes for k-flip neighborhood //NeighborhoodSize element of EOT
* @param _Kflip the number of flip to do EOT device_tmp;
*/ //Vector of neighborhoodSize copy of solution
T * vect;
void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping, bool mutex_kswap;
unsigned _Kflip) { bool mutex;
// the solution size
unsigned _size = _sol.size();
// Get Current solution fitness
Fitness fitness = _sol.fitness();
if (!mutex) {
//Allocate the space for solution in the device global memory
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
mutex = true;
}
//Copy the solution vector from the host to device
cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
cudaMemcpyHostToDevice);
//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);
//Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
}
protected:
IncrementEval & incrEval;
//NeighborhoodSize copy of solution
EOT device_setSolution;
//NeighborhoodSize element of EOT
EOT device_tmp;
//Vector of neighborhoodSize copy of solution
T * vect;
bool mutex_kswap;
bool mutex;
}; };
#endif #endif