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

Initial import.

Browse source 
git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@1986 331e1502-861f-0410-8da2-ba01fb791d7f
This commit is contained in:
boufaras 2010-11-03 10:21:45 +00:00
commit f1f07fe2a8
48 changed files with 6126 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

128
ParadisEO-GPU/src/eval/moCudaEval.h Normal file
View file

@ -0,0 +1,128 @@
/*
<moCudaEval.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Karima Boufaras, Thé Van LUONG
This software is governed by the CeCILL license under French law and
abiding by the rules of distribution of free software. You can use,
modify and/ or redistribute the software under the terms of the CeCILL
license as circulated by CEA, CNRS and INRIA at the following URL
"http://www.cecill.info".
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
with a limited warranty and the software's author, the holder of the
economic rights, and the successive licensors have only limited liability.
In this respect, the user's attention is drawn to the risks associated
with loading, using, modifying and/or developing or reproducing the
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
therefore means that it is reserved for developers and experienced
professionals having in-depth computer knowledge. Users are therefore
encouraged to load and test the software's suitability as regards their
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
same conditions as regards security.
The fact that you are presently reading this means that you have had
knowledge of the CeCILL license and that you accept its terms.
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
Contact: paradiseo-help@lists.gforge.inria.fr
*/
#ifndef moCudaEval_H
#define moCudaEval_H
#include <eval/moEval.h>
/**
* Abstract class for evaluation on GPU
*/
template<class Neighbor>
class moCudaEval: public moEval<Neighbor> {
public:
/**
* Define type of a solution corresponding to Neighbor
*/
typedef typename Neighbor::EOT EOT;
typedef typename EOT::Fitness Fitness;
/**
* Constructor
* @param
* _neighborhoodSize the size of the neighborhood
*/
moCudaEval(unsigned int _neighborhoodSize) {
neighborhoodSize = _neighborhoodSize;
host_FitnessArray = new Fitness[neighborhoodSize];
cudaMalloc((void**) &device_FitnessArray, neighborhoodSize
* sizeof(Fitness));
kernel_Dim=neighborhoodSize/BLOCK_SIZE+((neighborhoodSize%BLOCK_SIZE==0)?0:1);
}
/**
* Destructor
*/
~moCudaEval() {
delete[] host_FitnessArray;
cudaFree(device_FitnessArray);
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) {
_neighbor.fitness(host_FitnessArray[_neighbor.index()]);
/*std::cout << _sol.fitness() << " -host_FitnessArray["
<< _neighbor.index() << "]= "
<< host_FitnessArray[_neighbor.index()] << std::endl;*/
}
/**
* Compute fitness for all solution neighbors in device
* @param
* _sol the solution which generate the neighborhood
*/
virtual void neighborhoodEval(EOT & _sol, unsigned * _mapping,
unsigned _Kswap) {
}
/**
* Compute fitness for all solution neighbors in device
* @param
* _sol the solution which generate the neighborhood
*/
virtual void neighborhoodEval(EOT & _sol)=0;
protected:
//the host array to save all neighbors fitness
Fitness * host_FitnessArray;
//the device array to save neighbors fitness computed in device
Fitness * device_FitnessArray;
//the device solution
EOT device_solution;
//the size of neighborhood
unsigned int neighborhoodSize;
//Cuda kernel dimension
int kernel_Dim;
};
#endif

78
ParadisEO-GPU/src/eval/moCudaEvalFunc.h Normal file
View file

@ -0,0 +1,78 @@
/*
<moCudaEvalFunc.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Thé Van LUONG, Karima Boufaras
This software is governed by the CeCILL license under French law and
abiding by the rules of distribution of free software. You can use,
modify and/ or redistribute the software under the terms of the CeCILL
license as circulated by CEA, CNRS and INRIA at the following URL
"http://www.cecill.info".
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
with a limited warranty and the software's author, the holder of the
economic rights, and the successive licensors have only limited liability.
In this respect, the user's attention is drawn to the risks associated
with loading, using, modifying and/or developing or reproducing the
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
therefore means that it is reserved for developers and experienced
professionals having in-depth computer knowledge. Users are therefore
encouraged to load and test the software's suitability as regards their
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
same conditions as regards security.
The fact that you are presently reading this means that you have had
knowledge of the CeCILL license and that you accept its terms.
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
Contact: paradiseo-help@lists.gforge.inria.fr
*/
#ifndef __moCudaEvalFunc_H
#define __moCudaEvalFunc_H
/**
* Abstract class for Incremental evaluation of neighbor
*/
template<class Neighbor>
class moCudaEvalFunc {
public:
/**
* Define type of a solution corresponding to Neighbor
*/
typedef typename Neighbor::EOT EOT;
typedef typename EOT::Fitness Fitness;
/**
* Constructor
*/
moCudaEvalFunc() {
}
/**
* Destructor
*/
~moCudaEvalFunc() {
}
/**
*Virtual functor to compute fitness of a solution neighbor
*@param _solution the solution which generate the neighborhood
*@param _fitness the current solution fitness
*@param _id the index neighbor
*/
virtual inline __host__ __device__ Fitness operator() (EOT & _solution,Fitness _fitness, unsigned _id)=0;
};
#endif

121
ParadisEO-GPU/src/eval/moCudaKswapEval.h Normal file
View file

@ -0,0 +1,121 @@
/*
<moCudaKswapEval.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Karima Boufaras, Thé Van LUONG
This software is governed by the CeCILL license under French law and
abiding by the rules of distribution of free software. You can use,
modify and/ or redistribute the software under the terms of the CeCILL
license as circulated by CEA, CNRS and INRIA at the following URL
"http://www.cecill.info".
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
with a limited warranty and the software's author, the holder of the
economic rights, and the successive licensors have only limited liability.
In this respect, the user's attention is drawn to the risks associated
with loading, using, modifying and/or developing or reproducing the
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
therefore means that it is reserved for developers and experienced
professionals having in-depth computer knowledge. Users are therefore
encouraged to load and test the software's suitability as regards their
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
same conditions as regards security.
The fact that you are presently reading this means that you have had
knowledge of the CeCILL license and that you accept its terms.
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
Contact: paradiseo-help@lists.gforge.inria.fr
*/
#ifndef moCudaKswapEval_H
#define moCudaKswapEval_H
#include <eval/moCudaEval.h>
#include <eval/moCudakernelEval.h>
/**
* class for the cuda evaluation
*/
template<class Neighbor, class IncrementEval>
class moCudaKswapEval: public moCudaEval<Neighbor> {
public:
/**
* Define type of a solution corresponding to Neighbor
*/
typedef typename Neighbor::EOT EOT;
/**
* Define type of a vector corresponding to Solution
*/
typedef typename EOT::ElemType T;
/**
* Define type of a fitness corresponding to Solution
*/
typedef typename EOT::Fitness Fitness;
using moCudaEval<Neighbor>::neighborhoodSize;
using moCudaEval<Neighbor>::host_FitnessArray;
using moCudaEval<Neighbor>::device_FitnessArray;
using moCudaEval<Neighbor>::device_solution;
using moCudaEval<Neighbor>::kernel_Dim;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
* @param _incrEval the incremental evaluation
*/
moCudaKswapEval(unsigned _neighborhoodSize, IncrementEval & _incrEval) :
moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
}
/**
* Compute fitness for all solution neighbors in device
* @param _sol the solution which generate the neighborhood
*/
virtual void neighborhoodEval(EOT & _sol) {
}
/**
* Compute fitness for all solution neighbors in device
* @param _sol the solution which generate the neighborhood
* @param _mapping the array of indices mapping
* @param _Kswap the number of swap
*/
void neighborhoodEval(EOT & _sol, unsigned * _mapping, unsigned _Kswap) {
// the solution vector size
unsigned _size = _sol.size();
// Get Current solution fitness
Fitness fitness = _sol.fitness();
//Allocate the space for solution in the global memory of device
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
//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 neighbors fitness
kernelKswap<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap);
//Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
}
protected:
IncrementEval & incrEval;
};
#endif

113
ParadisEO-GPU/src/eval/moCudaVectorEval.h Normal file
View file

@ -0,0 +1,113 @@
/*
<moCudaVectorEval.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Karima Boufaras, Thé Van LUONG
This software is governed by the CeCILL license under French law and
abiding by the rules of distribution of free software. You can use,
modify and/ or redistribute the software under the terms of the CeCILL
license as circulated by CEA, CNRS and INRIA at the following URL
"http://www.cecill.info".
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
with a limited warranty and the software's author, the holder of the
economic rights, and the successive licensors have only limited liability.
In this respect, the user's attention is drawn to the risks associated
with loading, using, modifying and/or developing or reproducing the
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
therefore means that it is reserved for developers and experienced
professionals having in-depth computer knowledge. Users are therefore
encouraged to load and test the software's suitability as regards their
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
same conditions as regards security.
The fact that you are presently reading this means that you have had
knowledge of the CeCILL license and that you accept its terms.
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
Contact: paradiseo-help@lists.gforge.inria.fr
*/
#ifndef moCudaVectorEval_H
#define moCudavectorEval_H
#include <eval/moCudaEval.h>
#include <eval/moCudakernelEval.h>
/**
* class for the cuda evaluation
*/
template<class Neighbor, class IncrementEval>
class moCudaVectorEval: public moCudaEval<Neighbor> {
public:
/**
* Define type of a solution corresponding to Neighbor
*/
typedef typename Neighbor::EOT EOT;
/**
* Define type of a vector corresponding to Solution
*/
typedef typename EOT::ElemType T;
/**
* Define type of a fitness corresponding to Solution
*/
typedef typename EOT::Fitness Fitness;
using moCudaEval<Neighbor>::neighborhoodSize;
using moCudaEval<Neighbor>::host_FitnessArray;
using moCudaEval<Neighbor>::device_FitnessArray;
using moCudaEval<Neighbor>::device_solution;
using moCudaEval<Neighbor>::kernel_Dim;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
* @param _incrEval the incremental evaluation
*/
moCudaVectorEval(unsigned _neighborhoodSize, IncrementEval & _incrEval) :
moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
}
/**
* Compute fitness for all solution neighbors in device
* @param _sol the solution which generate the neighborhood
*/
virtual void neighborhoodEval(EOT & _sol) {
// the solution vector size
unsigned _size = _sol.size();
// Get Current solution fitness
Fitness fitness = _sol.fitness();
//Allocate the space for solution in the global memory of device
cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
//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 neighbors fitness
kernelEval<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize);
//Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
}
protected:
IncrementEval & incrEval;
};
#endif

96
ParadisEO-GPU/src/eval/moCudakernelEval.h Normal file
View file

@ -0,0 +1,96 @@
/*
<moCudakernelEval.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Karima Boufaras, Thé Van LUONG
This software is governed by the CeCILL license under French law and
abiding by the rules of distribution of free software. You can use,
modify and/ or redistribute the software under the terms of the CeCILL
license as circulated by CEA, CNRS and INRIA at the following URL
"http://www.cecill.info".
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
with a limited warranty and the software's author, the holder of the
economic rights, and the successive licensors have only limited liability.
In this respect, the user's attention is drawn to the risks associated
with loading, using, modifying and/or developing or reproducing the
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
therefore means that it is reserved for developers and experienced
professionals having in-depth computer knowledge. Users are therefore
encouraged to load and test the software's suitability as regards their
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
same conditions as regards security.
The fact that you are presently reading this means that you have had
knowledge of the CeCILL license and that you accept its terms.
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
Contact: paradiseo-help@lists.gforge.inria.fr
*/
#ifndef __moCudakernelEval_H
#define __moCudakernelEval_H
/**
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
* @param _eval how to evaluate each neighbor
* @param _solution representation of one max problem solution
* @param _allFitness Array of Fitness type to save all neighbors fitness
* @param _fitness the current solution fitness
* @param _neighborhoodsize the size of the neighborhood
*/
template<class EOT, class Fitness, class Neighbor, class IncrementEval>
__global__ void kernelEval(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
Fitness _fitness, unsigned _neighborhoodsize) {
// The thread identifier within a grid block's
int id = blockIdx.x * blockDim.x + threadIdx.x;
// In this representation each id identify one and only one neighbor in neighborhood
if (id < _neighborhoodsize) {
//Compute fitness for id'th neighbor
_allFitness[id] = _eval(_solution, _fitness, id);
}
}
/**
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
* @param _eval how to evaluate each neighbor
* @param _solution representation of one max problem solution
* @param _allFitness Array of Fitness type to save all neighbors fitness
* @param _fitness the current solution fitness
* @param _neighborhoodsize the size of the neighborhood
*/
template<class EOT, class Fitness, class Neighbor, class IncrementEval>
__global__ void kernelKswap(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _Kswap) {
// The thread identifier within a grid block's
int id = blockIdx.x * blockDim.x + threadIdx.x;
unsigned tmp;
unsigned i;
unsigned index;
// In this representation each id identify one and only one neighbor in neighborhood
if (id < _neighborhoodsize) {
tmp=_fitness;
for(i=0;i<=_Kswap;i++){
index=_mapping[id + i * _neighborhoodsize];
tmp= _eval(_solution, tmp, index);
}
_allFitness[id]=tmp;
}
}
#endif