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@2199 331e1502-861f-0410-8da2-ba01fb791d7f

Browse source
This commit is contained in:
boufaras 2011-03-17 13:46:57 +00:00
commit 43e1f0614f
1 changed files with 91 additions and 88 deletions
Download patch file Download diff file Expand all files Collapse all files

139
branches/ParadisEO-GPU/src/eval/moCudaEval.h
View file

@ -39,97 +39,100 @@
/** /**
* Abstract class for evaluation on GPU * Abstract class for evaluation on GPU
*/ */
template<class Neighbor> template<class Neighbor>
class moCudaEval: public moEval<Neighbor> { class moCudaEval: public moEval<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;
typedef typename EOT::Fitness Fitness; typedef typename EOT::Fitness Fitness;
/** /**
* Constructor * Constructor
* @param _neighborhoodSize the size of the neighborhood * @param _neighborhoodSize the size of the neighborhood
*/ */
moCudaEval(unsigned int _neighborhoodSize) { moCudaEval(unsigned int _neighborhoodSize) {
neighborhoodSize = _neighborhoodSize; neighborhoodSize = _neighborhoodSize;
host_FitnessArray = new Fitness[neighborhoodSize]; host_FitnessArray = new Fitness[neighborhoodSize];
cudaMalloc((void**) &device_FitnessArray, neighborhoodSize cudaMalloc((void**) &device_FitnessArray, neighborhoodSize
* sizeof(Fitness)); * sizeof(Fitness));
kernel_Dim = neighborhoodSize / BLOCK_SIZE + ((neighborhoodSize kernel_Dim = neighborhoodSize / BLOCK_SIZE + ((neighborhoodSize
% BLOCK_SIZE == 0) ? 0 : 1); % BLOCK_SIZE == 0) ? 0 : 1);
} }
/** /**
* Destructor * Destructor
*/ */
~moCudaEval() { ~moCudaEval() {
delete[] host_FitnessArray; delete[] host_FitnessArray;
cudaFree(device_FitnessArray); cudaFree(device_FitnessArray);
cudaFree(&device_solution); cudaFree(&device_solution);
}
/** }
* Set fitness of a solution neighbors
*@param _sol the solution which generate the neighborhood
*@param _neighbor the current neighbor
*/
void operator()(EOT & _sol, Neighbor & _neighbor) { /**
* Set fitness of a solution neighbors
*@param _sol the solution which generate the neighborhood
*@param _neighbor the current neighbor
*/
_neighbor.fitness(host_FitnessArray[_neighbor.index()]); void operator()(EOT & _sol, Neighbor & _neighbor) {
} _neighbor.fitness(host_FitnessArray[_neighbor.index()]);
/** }
* Compute fitness for all Kswap solution neighbors in device
* @param _sol the solution which generate the neighborhood
* @param _mapping the neighborhood mapping
* @param _Kswap the number of swap
*/
virtual void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping,
unsigned _Kswap) {
} /**
* Compute fitness for all Kswap solution neighbors in device
* @param _sol the solution which generate the neighborhood
* @param _mapping the neighborhood mapping
* @param _Kswap the number of swap
*/
/** virtual void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping,
* Compute fitness for all Kflip solution neighbors in device unsigned _Kswap) {
* @param _sol the solution which generate the neighborhood
* @param _mapping the neighborhood mapping
* @param _Kflip the number of bit to flip
*/
virtual void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping,
unsigned _Kflip) {
} }
/** /**
* Compute fitness for all solution neighbors in device * Compute fitness for all Kflip solution neighbors in device
* @param _sol the solution which generate the neighborhood * @param _sol the solution which generate the neighborhood
*/ * @param _mapping the neighborhood mapping
* @param _Kflip the number of bit to flip
*/
virtual void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping,
unsigned _Kflip) {
virtual void neighborhoodEval(EOT & _sol)=0; }
protected: /**
* Compute fitness for all solution neighbors in device
* @param _sol the solution which generate the neighborhood
*/
//the host array to save all neighbors fitness virtual void neighborhoodEval(EOT & _sol)=0;
Fitness * host_FitnessArray;
//the device array to save neighbors fitness computed in device protected:
Fitness * device_FitnessArray;
//the device solution //the host array to save all neighbors fitness
EOT device_solution; Fitness * host_FitnessArray;
//the size of neighborhood //the device array to save neighbors fitness computed in device
unsigned int neighborhoodSize; Fitness * device_FitnessArray;
//Cuda kernel dimension //the device solution
int kernel_Dim; EOT device_solution;
//the size of neighborhood
unsigned int neighborhoodSize;
//Cuda kernel dimension
int kernel_Dim;
}; };