Update the package
git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2247 331e1502-861f-0410-8da2-ba01fb791d7f
This commit is contained in:
boufaras
parent
1bbabf088c
commit
8b2cd7004e
5 changed files with 0 additions and 754 deletions
|
|
@ -1,139 +0,0 @@
|
||||||
/*
|
|
||||||
<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()]);
|
|
||||||
|
|
||||||
}
|
|
||||||
|
|
||||||
/**
|
|
||||||
* 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 Kflip solution neighbors in device
|
|
||||||
* @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
|
|
||||||
* @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
|
|
||||||
|
|
@ -1,84 +0,0 @@
|
||||||
/*
|
|
||||||
<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 CUDA 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::ElemType T;
|
|
||||||
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 _index the index neighbor
|
|
||||||
*/
|
|
||||||
|
|
||||||
virtual inline __host__ __device__ Fitness operator() (EOT & _solution,Fitness _fitness, unsigned int * _index){
|
|
||||||
|
|
||||||
return _fitness;
|
|
||||||
|
|
||||||
};
|
|
||||||
|
|
||||||
|
|
||||||
};
|
|
||||||
#endif
|
|
||||||
|
|
@ -1,212 +0,0 @@
|
||||||
/*
|
|
||||||
<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 K-swap neighborhood 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 int _neighborhoodSize, IncrementEval & _incrEval) :
|
|
||||||
moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
|
|
||||||
mutex = 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
|
|
||||||
cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
|
|
||||||
cudaMemcpyHostToDevice);
|
|
||||||
|
|
||||||
//Launch the Kernel to compute all permutation neighbors fitness
|
|
||||||
kernelPermutation<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_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++) {
|
|
||||||
for (int j = 0; j < _size; j++) {
|
|
||||||
vect[j + i * _size] = _sol.vect[j];
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
//Copy the set of solution from the host to device
|
|
||||||
cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size
|
|
||||||
* sizeof(T), cudaMemcpyHostToDevice);
|
|
||||||
|
|
||||||
//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);
|
|
||||||
|
|
||||||
//Copy the result from device to host
|
|
||||||
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
|
|
||||||
* sizeof(Fitness), cudaMemcpyDeviceToHost);
|
|
||||||
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
/**
|
|
||||||
* Compute fitness for all solution neighbors(K-flip of binary solution) in device
|
|
||||||
* @param _sol the solution which generate the neighborhood
|
|
||||||
* @param _mapping the array of mapping indexes for k-flip neighborhood
|
|
||||||
* @param _Kflip the number of flip to do
|
|
||||||
*/
|
|
||||||
|
|
||||||
void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping,
|
|
||||||
unsigned _Kflip) {
|
|
||||||
|
|
||||||
// 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
|
|
||||||
|
|
@ -1,113 +0,0 @@
|
||||||
/*
|
|
||||||
<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 parallel evaluation of neighborhood without mapping
|
|
||||||
*/
|
|
||||||
|
|
||||||
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 int _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
|
|
||||||
|
|
||||||
|
|
@ -1,206 +0,0 @@
|
||||||
/*
|
|
||||||
<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
|
|
||||||
* with linear mapping
|
|
||||||
* @param _eval how to evaluate each neighbor
|
|
||||||
* @param _solution the representation of solution( vector of int,float....)
|
|
||||||
* @param _allFitness the array of Fitness 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;
|
|
||||||
// array to save index to be changed
|
|
||||||
unsigned int index[1];
|
|
||||||
// In this representation each id identify one and only one neighbor in neighborhood
|
|
||||||
if (id < _neighborhoodsize) {
|
|
||||||
//Change the id'th element of solution
|
|
||||||
index[0]=id;
|
|
||||||
//Compute fitness for id'th neighbor
|
|
||||||
_allFitness[id] = _eval(_solution, _fitness,index);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
||||||
|
|
||||||
/**
|
|
||||||
* The kernel function called from the host and executed in device to compute all flip neighbors fitness at one time
|
|
||||||
* @param _eval how to evaluate each neighbor
|
|
||||||
* @param _solution representation of solution to flip
|
|
||||||
* @param _allFitness Array of Fitness type to save all neighbors fitness
|
|
||||||
* @param _fitness the current solution fitness
|
|
||||||
* @param _neighborhoodsize the size of the neighborhood
|
|
||||||
* @param _mapping the neighborhood mapping
|
|
||||||
* @param _Kflip the number of bit to flip
|
|
||||||
*/
|
|
||||||
|
|
||||||
template<class EOT, class Fitness, class Neighbor, class IncrementEval>
|
|
||||||
|
|
||||||
__global__ void kernelKflip(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
|
|
||||||
Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _Kflip) {
|
|
||||||
|
|
||||||
// The thread identifier within a grid block's
|
|
||||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
|
||||||
//save temporary fitness
|
|
||||||
unsigned tmp_fitness;
|
|
||||||
//counter of number of flip to do
|
|
||||||
unsigned i;
|
|
||||||
// array to save index to be changed
|
|
||||||
unsigned index[1];
|
|
||||||
// In this representation each id identify one and only one neighbor in neighborhood
|
|
||||||
if (id < _neighborhoodsize) {
|
|
||||||
//Init fitness with fitness of solution
|
|
||||||
tmp_fitness=_fitness;
|
|
||||||
//Evaluate neighbor after Kflip
|
|
||||||
for(i=0;i<=_Kflip;i++){
|
|
||||||
//The designed index to flip
|
|
||||||
index[0]=_mapping[id + i * _neighborhoodsize];
|
|
||||||
//Evaluate the neighbor
|
|
||||||
tmp_fitness= _eval(_solution, tmp_fitness, index);
|
|
||||||
|
|
||||||
}
|
|
||||||
//The final fitness of the Id'th neighbor
|
|
||||||
_allFitness[id]=tmp_fitness;
|
|
||||||
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
||||||
|
|
||||||
/**
|
|
||||||
* The kernel function called from the host and executed in device to compute all swap neighbors fitness at one time
|
|
||||||
* @param _eval how to evaluate each neighbor
|
|
||||||
* @param _solution representation ofsolution to swap
|
|
||||||
* @param _sol_tmp to save temporary a solution element to swap
|
|
||||||
* @param _allFitness Array of Fitness type to save all neighbors fitness
|
|
||||||
* @param _fitness the current solution fitness
|
|
||||||
* @param _neighborhoodsize the size of the neighborhood
|
|
||||||
* @param _mapping the neighborhood mapping
|
|
||||||
* @param _Kswap the number of swap to do
|
|
||||||
* @param _size the solution size
|
|
||||||
*/
|
|
||||||
|
|
||||||
template<class EOT,class Fitness, class Neighbor, class IncrementEval>
|
|
||||||
|
|
||||||
__global__ void kernelKswap(IncrementEval _eval,EOT _solution ,EOT _sol_tmp, Fitness* _allFitness,
|
|
||||||
Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _Kswap,unsigned _size) {
|
|
||||||
|
|
||||||
// The thread identifier within a grid block's
|
|
||||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
|
||||||
//save temporary fitness
|
|
||||||
int tmp_fitness;
|
|
||||||
//counter of number of swap to do
|
|
||||||
unsigned i;
|
|
||||||
// array to save index to be changed, solution size & thread id
|
|
||||||
unsigned index[4];
|
|
||||||
// In this representation each id identify one and only one neighbor in neighborhood
|
|
||||||
if (id < _neighborhoodsize) {
|
|
||||||
//the first index to swap
|
|
||||||
index[0]=_mapping[id];
|
|
||||||
//the second index to swap
|
|
||||||
index[1]=_mapping[id +_neighborhoodsize];
|
|
||||||
//the solution size
|
|
||||||
index[2]=_size;
|
|
||||||
//the thread id
|
|
||||||
index[3]=id;
|
|
||||||
//Init the temporary fitness with the initial solution fitness
|
|
||||||
tmp_fitness=_fitness;
|
|
||||||
|
|
||||||
//Evaluate neighbor after K-swap
|
|
||||||
for(i=2;i<=_Kswap+1;i++){
|
|
||||||
//Evaluate neighbor with index case
|
|
||||||
tmp_fitness=_eval(_solution, tmp_fitness, index);
|
|
||||||
//Permut the solution
|
|
||||||
_sol_tmp[id]=_solution[index[0]+id*index[2]];
|
|
||||||
_solution[index[0]+id*index[2]]=_solution[index[1]+id*index[2]];
|
|
||||||
_solution[index[1]+id*index[2]]=_sol_tmp[id];
|
|
||||||
//Init the next swap to do
|
|
||||||
index[0]=index[1];
|
|
||||||
index[1]=_mapping[id +i*_neighborhoodsize];
|
|
||||||
|
|
||||||
}
|
|
||||||
//save the final fitness of the id'th neighbor
|
|
||||||
_allFitness[id]=tmp_fitness;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
||||||
|
|
||||||
/**
|
|
||||||
* The kernel function called from the host and executed in device to compute all permutation neighbors fitness at one time
|
|
||||||
* @param _eval how to evaluate each neighbor
|
|
||||||
* @param _solution representation of 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
|
|
||||||
* @param _mapping the neighborhood mapping
|
|
||||||
* @param _size the solution size
|
|
||||||
*/
|
|
||||||
|
|
||||||
template<class EOT, class Fitness, class Neighbor, class IncrementEval>
|
|
||||||
__global__ void kernelPermutation(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
|
|
||||||
Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _size) {
|
|
||||||
|
|
||||||
// The thread identifier within a grid block's
|
|
||||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
|
||||||
// array to save index to be changed, solution size
|
|
||||||
unsigned index[4];
|
|
||||||
// In this representation each id identify one and only one neighbor in neighborhood
|
|
||||||
if (id < _neighborhoodsize) {
|
|
||||||
//The first index of permutation
|
|
||||||
index[0]=_mapping[id];
|
|
||||||
//The second index of permutation
|
|
||||||
index[1]=_mapping[id +_neighborhoodsize];
|
|
||||||
//The solution size
|
|
||||||
index[2]=_size;
|
|
||||||
//Puch 0 in the 3 index
|
|
||||||
index[3]=0;
|
|
||||||
_allFitness[id]=_eval(_solution,_fitness,index);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
#endif
|
|
||||||
Loading…
Add table
Add a link
Reference in a new issue