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

Add for new Concepts : Calibration ,unification of methods

Browse source 
git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2244 331e1502-861f-0410-8da2-ba01fb791d7f
This commit is contained in:
boufaras 2011-05-25 21:25:34 +00:00
commit 89f469a3f4
4 changed files with 623 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

146
branches/ParadisEO-GPU/src/eval/moGPUEval.h Normal file
View file

@ -0,0 +1,146 @@
/*
<moGPUEval.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 moGPUEval_H
#define moGPUEval_H
#include <eval/moEval.h>
/**
* Abstract class for evaluation on GPU
*/
template<class Neighbor>
class moGPUEval: public moEval<Neighbor> {
public:
/**
* Define type of a solution corresponding to Neighbor
**/
typedef typename Neighbor::EOT EOT;
/**
* Define type of a fitness corresponding to Solution
**/
typedef typename EOT::Fitness Fitness;
/**
* Define type of a vector corresponding to Solution
*/
typedef typename EOT::ElemType T;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
*/
moGPUEval(unsigned int _neighborhoodSize) {
neighborhoodSize = _neighborhoodSize;
host_FitnessArray = new Fitness[neighborhoodSize];
cudaMalloc((void**) &device_FitnessArray, neighborhoodSize
* sizeof(Fitness));
mutex = false;
#ifdef BLOCK_SIZE
NEW_kernel_Dim = neighborhoodSize / BLOCK_SIZE + ((neighborhoodSize
% BLOCK_SIZE == 0) ? 0 : 1);
NEW_BLOCK_SIZE = BLOCK_SIZE;
#endif
}
/**
* Destructor
*/
~moGPUEval() {
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 solution neighbors in device
* @param _sol the solution which generate the neighborhood
* @param _cpySolution Launch kernel with local copy option of solution in each thread
* @param _withCalibration an automatic configuration of kernel to launch( NB_THREAD BY BLOCK & NB_BLOCK BY KERNEL ),default (1)
*/
virtual void neighborhoodEval(EOT & _sol, bool _cpySolution,
bool _withCalibration) {
}
/**
* Compute fitness for all solution neighbors in device
* @param _sol the solution which generate the neighborhood
* @param _mapping the associated neighborhood mapping
* @param _cpySolution Launch kernel with local copy option of solution in each thread
* @param _withCalibration an automatic configuration of kernel to launch( NB_THREAD BY BLOCK & NB_BLOCK BY KERNEL ), default (1)
*/
virtual void neighborhoodEval(EOT & _sol, unsigned int * _mapping,
bool _cpySolution, bool _withCalibration) {
}
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;
//Assure mapping transfer from CPU memory to GPU global memory for once time
unsigned mutex;
//GPU kernel dimension
unsigned NEW_BLOCK_SIZE;
unsigned NEW_kernel_Dim;
};
#endif

197
branches/ParadisEO-GPU/src/eval/moGPUEvalByCpy.h Normal file
View file

@ -0,0 +1,197 @@
/*
<moGPUEvalByCpy.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 __moGPUEvalByCpy_H
#define __moGPUEvalByCpy_H
#include <eval/moGPUEval.h>
#include <eval/moGPUKernelEvalByCpy.h>
#include <performance/moGPUTimer.h>
/**
* class for the parallel evaluation of neighborhood
*/
template<class Neighbor, class Eval>
class moGPUEvalByCpy: public moGPUEval<Neighbor> {
public:
/**
* Define type of a solution corresponding to Neighbor
**/
typedef typename Neighbor::EOT EOT;
/**
* Define vector type of vector corresponding to Solution
**/
typedef typename EOT::ElemType T;
/**
* Define type of a fitness corresponding to Solution
**/
typedef typename EOT::Fitness Fitness;
using moGPUEval<Neighbor>::neighborhoodSize;
using moGPUEval<Neighbor>::host_FitnessArray;
using moGPUEval<Neighbor>::device_FitnessArray;
using moGPUEval<Neighbor>::device_solution;
using moGPUEval<Neighbor>::NEW_kernel_Dim;
using moGPUEval<Neighbor>::NEW_BLOCK_SIZE;
using moGPUEval<Neighbor>::mutex;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
* @param _eval how to evaluate a neighbor
*/
moGPUEvalByCpy(unsigned int _neighborhoodSize, Eval & _eval) :
moGPUEval<Neighbor> (_neighborhoodSize), eval(_eval) {
}
/**
* Compute fitness for all solution neighbors in device
* @param _cpySolution Launch kernel with local copy option of solution in each thread
* @param _withCalibration an automatic configuration of kernel to launch( NB_THREAD BY BLOCK & NB_BLOCK BY KERNEL ),default (1)
*/
void neighborhoodEval(EOT & _sol, bool _cpySolution, bool _withCalibration) {
if (_cpySolution) {
unsigned size=_sol.size();
if (!mutex) {
//Allocate the space for solution in the global memory of device
cudaMalloc((void**) &device_solution.vect, size * sizeof(T));
if (_withCalibration)
calibration(_sol);
mutex = true;
}
// Get Current solution fitness
Fitness fitness = _sol.fitness();
//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 with local copy of solution
moGPUKernelEvalByCpy<T,Fitness,Eval><<<NEW_kernel_Dim,NEW_BLOCK_SIZE >>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
//Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
} else
cout << "It's evaluation by copy set cpySolution to true" << endl;
}
virtual void calibration(EOT & _sol) {
unsigned size = _sol.size();
Fitness fitness = _sol.fitness();
unsigned NB_THREAD[6] = { 16, 32, 64, 128, 256, 512 };
double mean_time[7] = { 0, 0, 0, 0, 0, 0 };
unsigned i = 0;
double best_time = RAND_MAX;
unsigned tmp_kernel_Dim;
#ifndef BLOCK_SIZE
do {
tmp_kernel_Dim = neighborhoodSize / NB_THREAD[i]
+ ((neighborhoodSize % NB_THREAD[i] == 0) ? 0 : 1);
for (unsigned k = 0; k < 5; k++) {
cudaMemcpy(device_solution.vect, _sol.vect, size * sizeof(T),
cudaMemcpyHostToDevice);
moGPUTimer timer;
timer.start();
moGPUKernelEvalByCpy<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i] >>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
timer.deleteTimer();
}
if (best_time >= (mean_time[i] / 5)) {
best_time = mean_time[i] / 5;
NEW_BLOCK_SIZE = NB_THREAD[i];
NEW_kernel_Dim = tmp_kernel_Dim;
}
i++;
} while (i < 6);
#else
tmp_kernel_Dim =NEW_kernel_Dim;
for (unsigned k = 0; k < 5; k++) {
cudaMemcpy(device_solution.vect, _sol.vect, size * sizeof(T),
cudaMemcpyHostToDevice);
moGPUTimer timer;
timer.start();
moGPUKernelEvalByCpy<T,Fitness,Eval><<<tmp_kernel_Dim,BLOCK_SIZE >>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
timer.stop();
mean_time[6] += (timer.getTime());
timer.deleteTimer();
}
if (best_time >= (mean_time[6] / 5))
best_time = mean_time[6] / 5;
do {
tmp_kernel_Dim = neighborhoodSize / NB_THREAD[i]
+ ((neighborhoodSize % NB_THREAD[i] == 0) ? 0 : 1);
for (unsigned k = 0; k < 5; k++) {
cudaMemcpy(device_solution.vect, _sol.vect, size * sizeof(T),
cudaMemcpyHostToDevice);
moGPUTimer timer;
timer.start();
moGPUKernelEvalByCpy<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i] >>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
timer.deleteTimer();
}
if (best_time >= (mean_time[i] / 5)) {
best_time = mean_time[i] / 5;
NEW_BLOCK_SIZE = NB_THREAD[i];
NEW_kernel_Dim = tmp_kernel_Dim;
}
i++;
}while (i < 6);
#endif
}
protected:
Eval & eval;
};
#endif

196
branches/ParadisEO-GPU/src/eval/moGPUEvalByModif.h Normal file
View file

@ -0,0 +1,196 @@
/*
<moGPUEvalByModif.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 __moGPUEvalByModif_H
#define __moGPUEvalByModif_H
#include <eval/moGPUKernelEvalByModif.h>
#include <eval/moGPUEval.h>
#include <performance/moGPUTimer.h>
/**
* class for the parallel evaluation of neighborhood
*/
template<class Neighbor, class Eval>
class moGPUEvalByModif: public moGPUEval<Neighbor> {
public:
/**
* Define type of a solution corresponding to Neighbor
**/
typedef typename Neighbor::EOT EOT;
/**
* Define vector type of vector corresponding to Solution
**/
typedef typename EOT::ElemType T;
/**
* Define type of a fitness corresponding to Solution
**/
typedef typename EOT::Fitness Fitness;
using moGPUEval<Neighbor>::neighborhoodSize;
using moGPUEval<Neighbor>::host_FitnessArray;
using moGPUEval<Neighbor>::device_FitnessArray;
using moGPUEval<Neighbor>::device_solution;
using moGPUEval<Neighbor>::NEW_kernel_Dim;
using moGPUEval<Neighbor>::NEW_BLOCK_SIZE;
using moGPUEval<Neighbor>::mutex;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
* @param _eval the incremental evaluation
*/
moGPUEvalByModif(unsigned int _neighborhoodSize, Eval & _eval) :
moGPUEval<Neighbor> (_neighborhoodSize), eval(_eval) {
}
/**
* Compute fitness for all solution neighbors in device
* @param _cpySolution Launch kernel with local copy option of solution in each thread
* @param _withCalibration an automatic configuration of kernel to launch( NB_THREAD BY BLOCK & NB_BLOCK BY KERNEL ),default (1)
*/
void neighborhoodEval(EOT & _sol, bool _cpySolution, bool _withCalibration) {
if (!_cpySolution) {
unsigned size = _sol.size();
if (!mutex) {
//Allocate the space for solution in the global memory of device
cudaMalloc((void**) &device_solution.vect, size * sizeof(T));
if (_withCalibration)
calibration(_sol);
mutex = true;
}
// Get Current solution fitness
Fitness fitness = _sol.fitness();
//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
moGPUKernelEvalByModif<T,Fitness,Eval><<<NEW_kernel_Dim,NEW_BLOCK_SIZE>>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
//Copy the result from device to host
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
} else
cout << "It's evaluation by Modif set cpySolution to false"
<< endl;
}
virtual void calibration(EOT & _sol) {
unsigned size = _sol.size();
Fitness fitness = _sol.fitness();
unsigned NB_THREAD[6] = { 16, 32, 64, 128, 256, 512 };
double mean_time[7] = { 0, 0, 0, 0, 0, 0 };
unsigned i = 0;
double best_time = RAND_MAX;
unsigned tmp_kernel_Dim;
#ifndef BLOCK_SIZE
do {
tmp_kernel_Dim = neighborhoodSize / NB_THREAD[i]
+ ((neighborhoodSize % NB_THREAD[i] == 0) ? 0 : 1);
for (unsigned k = 0; k < 5; k++) {
cudaMemcpy(device_solution.vect, _sol.vect, size * sizeof(T),
cudaMemcpyHostToDevice);
moGPUTimer timer;
timer.start();
moGPUKernelEvalByModif<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i]>>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
timer.deleteTimer();
}
if (best_time >= (mean_time[i] / 5)) {
best_time = mean_time[i] / 5;
NEW_BLOCK_SIZE = NB_THREAD[i];
NEW_kernel_Dim = tmp_kernel_Dim;
}
i++;
} while (i < 6);
#else
tmp_kernel_Dim =NEW_kernel_Dim;
for (unsigned k = 0; k < 5; k++) {
cudaMemcpy(device_solution.vect, _sol.vect, size * sizeof(T),
cudaMemcpyHostToDevice);
moGPUTimer timer;
timer.start();
moGPUKernelEvalByModif<T,Fitness,Eval><<<tmp_kernel_Dim,BLOCK_SIZE>>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
timer.stop();
mean_time[6] += (timer.getTime());
timer.deleteTimer();
}
if (best_time >= (mean_time[6] / 5))
best_time = mean_time[6] / 5;
do {
tmp_kernel_Dim = neighborhoodSize / NB_THREAD[i]
+ ((neighborhoodSize % NB_THREAD[i] == 0) ? 0 : 1);
for (unsigned k = 0; k < 5; k++) {
cudaMemcpy(device_solution.vect, _sol.vect, size * sizeof(T),
cudaMemcpyHostToDevice);
moGPUTimer timer;
timer.start();
moGPUKernelEvalByModif<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i]>>>(eval,device_solution.vect,device_FitnessArray,fitness,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
timer.deleteTimer();
}
if (best_time >= (mean_time[i] / 5)) {
best_time = mean_time[i] / 5;
NEW_BLOCK_SIZE = NB_THREAD[i];
NEW_kernel_Dim = tmp_kernel_Dim;
}
i++;
}while (i < 6);
#endif
}
protected:
Eval & eval;
};
#endif

84
branches/ParadisEO-GPU/src/eval/moGPUEvalFunc.h Normal file
View file

@ -0,0 +1,84 @@
/*
<moGPUEvalFunc.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 __moGPUEvalFunc_H
#define __moGPUEvalFunc_H
/**
* Abstract class for GPU evaluation of neighbor
*/
template<class Neighbor>
class moGPUEvalFunc {
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
*/
moGPUEvalFunc() {
}
/**
* Destructor
*/
virtual ~moGPUEvalFunc() {
}
/**
*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 an array that contains a set of indexes corresponding to the current thread identifier neighbor
* the last element of this array contains neighborhood size
*/
virtual inline __host__ __device__ Fitness operator() (T * _solution,Fitness _fitness, unsigned int * _index) {
return _fitness;
}
};
#endif