nojhan/paradiseo
Watch
1
0
Fork
You've already forked paradiseo
 0
Code Issues Pull requests Projects Releases Packages Wiki Activity

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

Browse source
This commit is contained in:
boufaras 2011-10-25 13:27:22 +00:00
commit 1f10f11ed6
49 changed files with 5536 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

141
trunk/paradiseo-gpu/src/GPUType/moGPUBitVector.h Normal file
View file

@ -0,0 +1,141 @@
/*
<moGPUBitVector.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 __moGPUBitVector_H_
#define __moGPUBitVector_H_
#include <GPUType/moGPUVector.h>
#include <stdlib.h>
/**
* Implementation of Bit vector representation on GPU.
*/
template<class Fitness>
class moGPUBitVector: public moGPUVector<bool, Fitness> {
public:
/**
* Define bool vector corresponding to Solution
**/
typedef bool ElemType;
using moGPUVector<bool, Fitness>::vect;
using moGPUVector<bool, Fitness>::N;
/**
* Default constructor.
*/
moGPUBitVector() :
moGPUVector<bool, Fitness> () {
}
/**
*Constructor.
*@param _neighborhoodSize The neighborhood size.
*/
moGPUBitVector(unsigned _neighborhoodSize) :
moGPUVector<bool, Fitness> (_neighborhoodSize) {
create();
}
/**
*Constructor.
*@param _neighborhoodSize The neighborhood size.
*@param _b Value to assign to vector.
*/
moGPUBitVector(unsigned _neighborhoodSize, bool _b) :
moGPUVector<bool, Fitness> (_neighborhoodSize) {
for (unsigned i = 0; i < _neighborhoodSize; i++)
vect[i] = _b;
}
/**
*Initializer of random bit vector.
*/
void create() {
for (unsigned i = 0; i < N; i++) {
vect[i] = (int) (rng.rand() / RAND_MAX);
}
}
/**
*Function inline to set the size of vector, called from host.
*@param _size the vector size
*/
void setSize(unsigned _size) {
if (_size < N) {
moGPUBitVector<Fitness> tmp_vect(_size);
for (unsigned i = 0; i < tmp_vect.N; i++)
tmp_vect.vect[i] = vect[i];
(tmp_vect).invalidate();
(*this) = tmp_vect;
} else if (_size > N) {
moGPUBitVector<Fitness> tmp_vect(_size);
for (unsigned i = 0; i < N; i++)
tmp_vect.vect[i] = vect[i];
(tmp_vect).invalidate();
(*this) = tmp_vect;
}
}
/**
* Write object. Called printOn since it prints the object _on_ a stream.
* @param _os A std::ostream.
*/
void printOn(std::ostream& _os) const {
EO<Fitness>::printOn(_os);
_os << ' ';
_os << N << ' ';
for (unsigned int i = 0; i < N; i++)
_os << (*this)[i] << ' ';
}
};
#endif

122
trunk/paradiseo-gpu/src/GPUType/moGPUIntVector.h Normal file
View file

@ -0,0 +1,122 @@
/*
<moGPUIntVector.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 __moGPUIntVector_H_
#define __moGPUIntVector_H_
#include <GPUType/moGPUVector.h>
#include <stdlib.h>
/**
* Implementation of integer vector representation on GPU.
*/
template<class Fitness>
class moGPUIntVector: public moGPUVector<int, Fitness> {
public:
using moGPUVector<int, Fitness>::vect;
using moGPUVector<int, Fitness>::N;
/**
* Default constructor.
*/
moGPUIntVector() :
moGPUVector<int, Fitness> () {
}
/**
*Constructor.
*@param _neighborhoodSize The neighborhood size.
*/
moGPUIntVector(unsigned _neighborhoodSize) :
moGPUVector<int, Fitness> (_neighborhoodSize) {
create();
}
/**
*Initializer of random integer vector.
*/
virtual void create() {
for (unsigned i = 0; i < N; i++)
vect[i] = ((int) (rng.rand())) % N + 1;
}
/**
*Function inline to set the size of vector, called from host.
*@param _size the vector size
*/
virtual inline __host__ void setSize(unsigned _size) {
if(N==_size)
return;
moGPUIntVector<Fitness> tmp_vect(_size);
if(_size<N) {
for (unsigned i = 0; i < tmp_vect.N; i++)
tmp_vect.vect[i]= vect[i];
}
else if(_size>N) {
for (unsigned i = 0; i <N; i++)
tmp_vect.vect[i]= vect[i];
}
(tmp_vect).invalidate();
(*this)=tmp_vect;
}
/**
* Write object. Called printOn since it prints the object _on_ a stream.
* @param _os A std::ostream.
*/
void printOn(std::ostream& _os) const {
EO<Fitness>::printOn(_os);
_os << ' ';
_os << N << ' ';
unsigned int i;
for (i = 0; i < N; i++)
_os << vect[i] << ' ';
}
};
#endif

92
trunk/paradiseo-gpu/src/GPUType/moGPUPermutationVector.h Normal file
View file

@ -0,0 +1,92 @@
/*
<moGPUPermutationVector.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 __moGPUPermutationVector_H_
#define __moGPUPermutationVector_H_
#include <GPUType/moGPUIntVector.h>
/**
* Implementation of Permutation vector representation on GPU.
*/
template<class Fitness>
class moGPUPermutationVector: public moGPUIntVector<Fitness> {
public:
using moGPUIntVector<Fitness>::vect;
using moGPUIntVector<Fitness>::N;
/**
* Default constructor.
*/
moGPUPermutationVector() :
moGPUIntVector<Fitness> () {
}
/**
*Constructor.
*@param _size The solution size.
*/
moGPUPermutationVector(unsigned _size) :
moGPUIntVector<Fitness> (_size) {
create();
}
/**
*Initializer of random permuatation vector.
*/
void create() {
unsigned random;
int tmp;
for (unsigned i = 0; i < N; i++)
vect[i] = i;
// we want a random permutation so we shuffle
for (unsigned i = 0; i < N; i++) {
random = rng.rand() % (N - i) + i;
tmp = vect[i];
vect[i] = vect[random];
vect[random] = tmp;
}
}
};
#endif

123
trunk/paradiseo-gpu/src/GPUType/moGPURealVector.h Normal file
View file

@ -0,0 +1,123 @@
/*
<moGPURealVector.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 __moGPURealVector_H_
#define __moGPURealVector_H_
#include <GPUType/moGPUVector.h>
#include <stdlib.h>
/**
* Implementation of real vector representation on GPU.
*/
template<class Fitness>
class moGPURealVector: public moGPUVector<float, Fitness> {
public:
using moGPUVector<float, Fitness>::vect;
using moGPUVector<float, Fitness>::N;
/**
* Default constructor.
*/
moGPURealVector() :
moGPUVector<float, Fitness> () {
}
/**
*Constructor.
*@param _neighborhoodSize The neighborhood size.
*/
moGPURealVector(unsigned _neighborhoodSize) :
moGPUVector<float, Fitness> (_neighborhoodSize) {
create();
}
/**
*Initializer of random real vector.
*/
void create() {
for (unsigned i = 0; i < N; i++)
vect[i] = (float) rng.rand() / RAND_MAX;
}
/**
*Function inline to set the size of vector, called from host.
*@param _size the vector size
*/
virtual inline __host__ void setSize(unsigned _size) {
if(_size<N) {
moGPURealVector<Fitness> tmp_vect(_size);
for (unsigned i = 0; i < tmp_vect.N; i++)
tmp_vect.vect[i]= vect[i];
(tmp_vect).invalidate();
(*this)=tmp_vect;
}
else if(_size>N) {
moGPURealVector<Fitness> tmp_vect(_size);
for (unsigned i = 0; i <N; i++)
tmp_vect.vect[i]= vect[i];
(tmp_vect).invalidate();
(*this)=tmp_vect;
}
}
/**
* Write object. Called printOn since it prints the object _on_ a stream.
* @param _os A std::ostream.
*/
void printOn(std::ostream& _os) const {
EO<Fitness>::printOn(_os);
_os << ' ';
_os << N << ' ';
unsigned int i;
for (i = 0; i < N; i++)
_os << vect[i] << ' ';
}
};
#endif

184
trunk/paradiseo-gpu/src/GPUType/moGPUVector.h Normal file
View file

@ -0,0 +1,184 @@
/*
<moGPUVector.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 __moGPUVector_H_
#define __moGPUVector_H_
#include <eo>
/**
* Implementation of a GPU solution representation.
*/
template<class ElemT, class Fitness>
class moGPUVector: public EO<Fitness> {
public:
/**
* Define vector type corresponding to Solution
*/
typedef ElemT ElemType;
/**
* Default constructor.
*/
moGPUVector() :
N(0) {
}
/**
*Constructor.
*@param _size The solution size.
*/
moGPUVector(unsigned _size) :
N(_size) {
vect = new ElemType[N];
}
/**
*Copy Constructor
*@param _vector The vector passed to the function to determine the new content.
*/
moGPUVector(const moGPUVector & _vector) {
N = _vector.N;
vect = new ElemType[N];
for (unsigned i = 0; i < N; i++)
vect[i] = _vector.vect[i];
if (!(_vector.invalid()))
fitness(_vector.fitness());
else
(*this).invalidate();
}
/**
* Destructor.
*/
~moGPUVector() {
if (N >= 1)
delete[] vect;
}
/**
*How to fill the solution vector.
*/
virtual void create() =0;
/**
*Assignment operator
*@param _vector The vector passed to the function to determine the new content.
*@return a new vector.
*/
moGPUVector& operator=(const moGPUVector & _vector) {
if (!(N == _vector.N)) {
N = _vector.N;
vect = new ElemType[N];
}
for (unsigned i = 0; i < N; i++){
vect[i] = _vector[i];
}
if (!(_vector.invalid()))
fitness(_vector.fitness());
else
(*this).invalidate();
return (*this);
}
/**
*An accessor read only on the i'th element of the vector (function inline can be called from host or device).
*@param _i The i'th element of vector.
*@return The i'th element of the vector for read only
*/
inline __host__ __device__ const ElemType & operator[](unsigned _i) const {
return vect[_i];
}
/**
*An accessor read-write on the i'th element of the vector(function inline can be called from host or device).
*@param _i The i'th element of the vector.
*@return The i'th element of the vector for read-write
*/
inline __host__ __device__ ElemType & operator[](unsigned _i) {
return vect[_i];
}
/**
*Function inline to get the size of vector, called from host and device.
*@return The vector size's
*/
inline __host__ __device__ unsigned size() {
return N;
}
/**
* method to set the size of vector
*@param _size the vector size
*/
virtual void setSize(unsigned _size)=0;
/**
* Write object. Called printOn since it prints the object _on_ a stream.
* @param _os A std::ostream.
*/
virtual void printOn(std::ostream& _os) const=0;
protected:
ElemType * vect;
unsigned N;
};
#endif

146
trunk/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

200
trunk/paradiseo-gpu/src/eval/moGPUEvalByCpy.h Normal file
View file

@ -0,0 +1,200 @@
/*
<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 _sol the solution that generate the neighborhood to evaluate parallely
* @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;
}
/**
* Compute the best combination of number of block by grid and number of thread within block
* @param _sol the solution that generate the neighborhood to evaluate parallely
*/
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());
}
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());
}
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());
}
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

199
trunk/paradiseo-gpu/src/eval/moGPUEvalByModif.h Normal file
View file

@ -0,0 +1,199 @@
/*
<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 _sol the solution that generate the neighborhood to evaluate parallely
* @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;
}
/**
* Compute the best combination of number of block by grid and number of thread within block
* @param _sol the solution that generate the neighborhood to evaluate parallely
*/
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());
}
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());
}
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());
}
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
trunk/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

71
trunk/paradiseo-gpu/src/eval/moGPUKernelEvalByCpy.h Normal file
View file

@ -0,0 +1,71 @@
/*
<moGPUKernelEvalByCpy.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 __moGPUKernelEvalByCpy_H
#define __moGPUKernelEvalByCpy_H
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
/**
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
* without mapping, each thread id compute one fitness by copying localy the solution
* @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 T, class Fitness, class Eval>
__global__ void moGPUKernelEvalByCpy(Eval _eval, T * _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 index[2];
T sol_tmp[SIZE];
// In this representation each id identify one and only one neighbor in neighborhood
if (id < _neighborhoodsize) {
for(unsigned i=0;i<SIZE;i++)
sol_tmp[i]=_solution[i];
//Change the id'th element of solution
index[0]=id;
index[1]=_neighborhoodsize;
//Compute fitness for id'th neighbor
_allFitness[id] = _eval(sol_tmp,_fitness,index);
}
}
#endif

68
trunk/paradiseo-gpu/src/eval/moGPUKernelEvalByModif.h Normal file
View file

@ -0,0 +1,68 @@
/*
<moGPUKernelEvalByModif.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 __moGPUKernelEvalByModif_H
#define __moGPUKernelEvalByModif_H
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
/**
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
* without mapping, each thread id compute one fitness by modif of solution
* @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 T, class Fitness, class Eval>
__global__ void moGPUKernelEvalByModif(Eval _eval, T * _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 index[2];
// 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;
index[1]=_neighborhoodsize;
//Compute fitness for id'th neighbor
_allFitness[id] = _eval(_solution,_fitness,index);
}
}
#endif

202
trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByCpy.h Normal file
View file

@ -0,0 +1,202 @@
/*
<moGPUMappingEvalByCpy.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 __moGPUMappingEvalByCpy_H
#define __moGPUMappingEvalByCpy_H
#include <eval/moGPUEval.h>
#include <eval/moGPUMappingKernelEvalByCpy.h>
#include <performance/moGPUTimer.h>
/**
* class for the Mapping neighborhood evaluation
*/
template<class Neighbor, class Eval>
class moGPUMappingEvalByCpy: public moGPUEval<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 moGPUEval<Neighbor>::neighborhoodSize;
using moGPUEval<Neighbor>::host_FitnessArray;
using moGPUEval<Neighbor>::device_FitnessArray;
using moGPUEval<Neighbor>::device_solution;
using moGPUEval<Neighbor>::NEW_BLOCK_SIZE;
using moGPUEval<Neighbor>::NEW_kernel_Dim;
using moGPUEval<Neighbor>::mutex;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
* @param _eval how to evaluate a neighbor
*/
moGPUMappingEvalByCpy(unsigned int _neighborhoodSize, Eval & _eval) :
moGPUEval<Neighbor> (_neighborhoodSize), eval(_eval) {
}
/**
* Destructor
*/
~moGPUMappingEvalByCpy() {
}
/**
* Compute fitness for all solution neighbors in device with associated mapping
* @param _sol the solution that generate the neighborhood to evaluate parallely
* @param _mapping the array of mapping indexes that associate a neighbor identifier to X-position
* @param _cpySolution Launch kernel with local copy option of solution in each thread if it's set to true
* @param _withCalibration an automatic kernel configuration, fix nbr of thread by block and nbr of grid by kernel
*/
void neighborhoodEval(EOT & _sol, unsigned int * _mapping,
bool _cpySolution, bool _withCalibration) {
if (_cpySolution) {
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));
if (_withCalibration)
calibration(_sol, _mapping);
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 neighbors fitness,using a given mapping
moGPUMappingKernelEvalByCpy<T,Fitness,Eval><<<NEW_kernel_Dim,NEW_BLOCK_SIZE >>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
} else
cout << "It's evaluation by copy set cpySolution to true" << endl;
}
/**
* Compute the best combination of number of block by grid and number of thread within block
* @param _sol the solution that generate the neighborhood to evaluate parallely
* @param _mapping the array of mapping indexes that associate a neighbor identifier to X-position
*/
virtual void calibration(EOT & _sol, unsigned int * _mapping) {
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 = 0;
unsigned tmp_kernel_Dim;
best_time = RAND_MAX;
#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();
moGPUMappingKernelEvalByCpy<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i]>>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
}
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();
moGPUMappingKernelEvalByCpy<T,Fitness,Eval><<<tmp_kernel_Dim,BLOCK_SIZE>>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
timer.stop();
mean_time[6] += (timer.getTime());
}
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();
moGPUMappingKernelEvalByCpy<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i]>>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
}
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

200
trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByModif.h Normal file
View file

@ -0,0 +1,200 @@
/*
<moGPUMappingEvalByModif.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 __moGPUMappingEvalByModif_H
#define __moGPUMappingEvalByModif_H
#include <eval/moGPUEval.h>
#include <eval/moGPUMappingKernelEvalByModif.h>
#include <performance/moGPUTimer.h>
/**
* class for the Mapping neighborhood evaluation
*/
template<class Neighbor, class Eval>
class moGPUMappingEvalByModif: public moGPUEval<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 moGPUEval<Neighbor>::neighborhoodSize;
using moGPUEval<Neighbor>::host_FitnessArray;
using moGPUEval<Neighbor>::device_FitnessArray;
using moGPUEval<Neighbor>::device_solution;
using moGPUEval<Neighbor>::NEW_BLOCK_SIZE;
using moGPUEval<Neighbor>::NEW_kernel_Dim;
using moGPUEval<Neighbor>::mutex;
/**
* Constructor
* @param _neighborhoodSize the size of the neighborhood
* @param _eval the incremental evaluation
*/
moGPUMappingEvalByModif(unsigned int _neighborhoodSize, Eval & _eval) :
moGPUEval<Neighbor> (_neighborhoodSize), eval(_eval) {
}
/**
* Destructor
*/
~moGPUMappingEvalByModif() {
}
/**
* Compute fitness for all solution neighbors in device with associated mapping
* @param _sol the solution that generate the neighborhood to evaluate parallely
* @param _mapping the array of mapping indexes that associate a neighbor identifier to X-position
* @param _cpySolution Launch kernel with local copy option of solution in each thread if it's set to true
* @param _withCalibration an automatic kernel configuration, fix nbr of thread by block and nbr of grid by kernel
*/
void neighborhoodEval(EOT & _sol, unsigned int * _mapping,
bool _cpySolution, bool _withCalibration) {
if (!_cpySolution) {
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));
if (_withCalibration)
calibration(_sol, _mapping);
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 neighbors fitness,using a given mapping
moGPUMappingKernelEvalByModif<T,Fitness,Eval><<<NEW_kernel_Dim,NEW_BLOCK_SIZE >>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
* sizeof(Fitness), cudaMemcpyDeviceToHost);
} else
cout << "It's evaluation by Modif set cpySolution to false" << endl;
}
/**
* Compute the best combination of number of block by grid and number of thread within block
* @param _sol the solution that generate the neighborhood to evaluate parallely
* @param _mapping the array of mapping indexes that associate a neighbor identifier to X-position
*/
virtual void calibration(EOT & _sol, unsigned int * _mapping) {
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 = 0;
unsigned tmp_kernel_Dim;
best_time = RAND_MAX;
#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();
moGPUMappingKernelEvalByModif<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i]>>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
}
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();
moGPUMappingKernelEvalByModif<T,Fitness,Eval><<<tmp_kernel_Dim,NEW_BLOCK_SIZE >>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
timer.stop();
mean_time[6] += (timer.getTime());
}
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();
moGPUMappingKernelEvalByModif<T,Fitness,Eval><<<tmp_kernel_Dim,NB_THREAD[i] >>>(eval,device_solution.vect,device_FitnessArray,fitness,_mapping,neighborhoodSize);
timer.stop();
mean_time[i] += (timer.getTime());
}
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

76
trunk/paradiseo-gpu/src/eval/moGPUMappingKernelEvalByCpy.h Normal file
View file

@ -0,0 +1,76 @@
/*
<moGPUMappingKernelEvalByCpy.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 __moGPUMappingKernelEvalByCpy_H
#define __moGPUMappingKernelEvalByCpy_H
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
/**
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
* without mapping, each thread id compute one fitness by modif of solution
* @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 _mapping associate to each threadID a set of correspondent indexes
* @param _neighborhoodsize the size of the neighborhood
*/
template<class T, class Fitness, class Eval>
__global__ void moGPUMappingKernelEvalByCpy(Eval _eval, T * _solution, Fitness* _allFitness,
Fitness _fitness,unsigned * _mapping,unsigned _neighborhoodsize) {
// The thread identifier within a grid block's
int id = blockIdx.x * blockDim.x + threadIdx.x;
//counter of number of x-change
unsigned i;
// array to save set a set of indexes corresponding to the current thread identifier
unsigned index[NB_POS+2];
T sol_tmp[SIZE];
// In this representation each id identify one and only one neighbor in neighborhood
if (id < _neighborhoodsize) {
for(i=0;i<SIZE;i++)
sol_tmp[i]=_solution[i];
for(i=0;i<NB_POS;i++)
index[i]=_mapping[id + i * _neighborhoodsize];
index[NB_POS]=_neighborhoodsize;
index[NB_POS+1]=id;
//Evaluate by Modif Id'th neighbor with index mapping
_allFitness[id]=_eval(sol_tmp,_fitness, index);
}
}
#endif

73
trunk/paradiseo-gpu/src/eval/moGPUMappingKernelEvalByModif.h Normal file
View file

@ -0,0 +1,73 @@
/*
<moGPUMappingKernelEvalByModif.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 __moGPUMappingKernelEvalByModif_H
#define __moGPUMappingKernelEvalByModif_H
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
/**
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
* without mapping, each thread id compute one fitness by modif of solution
* @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 _mapping associate to each threadID a set of correspondent indexes
* @param _neighborhoodsize the size of the neighborhood
*/
template<class T, class Fitness, class Eval>
__global__ void moGPUMappingKernelEvalByModif(Eval _eval, T * _solution, Fitness* _allFitness,
Fitness _fitness,unsigned int * _mapping,unsigned _neighborhoodsize) {
// The thread identifier within a grid block's
int id = blockIdx.x * blockDim.x + threadIdx.x;
//counter of number of x-change
unsigned i;
// array to save set a set of indexes corresponding to the current thread identifier
unsigned index[NB_POS+2];
// In this representation each id identify one and only one neighbor in neighborhood
if (id < _neighborhoodsize) {
for(i=0;i<NB_POS;i++)
index[i]=_mapping[id + i * _neighborhoodsize];
index[NB_POS]=_neighborhoodsize;
index[NB_POS+1]=id;
//Evaluate by Modif Id'th neighbor with index mapping
_allFitness[id]=_eval(_solution,_fitness, index);
}
}
#endif

75
trunk/paradiseo-gpu/src/memory/moGPUAllocator.h Normal file
View file

@ -0,0 +1,75 @@
/*
<moGPUAllocator.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Boufaras Karima, 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 __moGPUAllocator_H_
#define __moGPUAllocator_H_
/**
* class for allocation data on GPU global memory
*/
class moGPUAllocator {
public:
/**
* Constructor
*/
moGPUAllocator() {
}
/**
*Allocate data on GPU global memory
*@param _data the data to allocate on GPU global memory
*@param _dataSize the size of data to allocate on GPU memory
*/
template<typename T>
void operator()(T* & _data, unsigned _dataSize) {
//Allocate data in GPU memory
cudaMalloc((void**) &_data, _dataSize * sizeof(T));
}
/**
* Destructor
*/
~moGPUAllocator() {
}
};
#endif

112
trunk/paradiseo-gpu/src/memory/moGPUCopy.h Normal file
View file

@ -0,0 +1,112 @@
/*
<moGPUCopy.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Boufaras Karima, 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 __moGPUCopy_H_
#define __moGPUCopy_H_
/**
* class to copy data from CPU memory to GPU global memory and vice versa
*/
class moGPUCopy {
public:
/**
* Constructor
*/
moGPUCopy() {
}
/**
*Copy data from CPU memory to GPU global memory (default copy)
*@param _data the data representation where the data will be copied
*@param _dataTocpy the data to copy from CPU memory to GPU memory
*@param _dataSize the size of data to copy
*/
template<typename T>
void operator()(T* & _data, T * & _dataTocpy, unsigned _dataSize) {
//copy data from CPU memory to GPU memory
cudaMemcpy(_data, _dataTocpy, _dataSize * sizeof(T),
cudaMemcpyHostToDevice);
}
/**
*Copy device data from GPU global memory to global variable declared in device
*@param _dev_data the device global variable
*@param _dataTocpy the data to copy GPU global memory to GPU global variable
*/
template<typename T>
void operator()(T* & _dev_data, T * & _dataTocpy) {
//Copy n bytes from the memory area pointed to by _dataTocpy to the memory area pointed to by offset bytes from the start of symbol _dev_data
cudaMemcpyToSymbol(_dev_data, &_dataTocpy, sizeof(_dataTocpy));
}
/**
*Copy data from CPU memory to GPU global memory and vice versa
*@param _data the data representation where the data will be copied
*@param _dataTocpy the data to copy from CPU memory to GPU memory and vice versa
*@param _dataSize the size of data to copy
*@param _HostToDevice the direction of copy(true if copy will be done from CPU memory to GPU memory)
*/
template<typename T>
void operator()(T* & _data, T * & _dataTocpy, unsigned _dataSize,
bool _HostToDevice) {
if (_HostToDevice) {
//copy data from CPU memory to GPU global memory
cudaMemcpy(_data, _dataTocpy, _dataSize * sizeof(T),
cudaMemcpyHostToDevice);
}
else {
//copy data from GPU global memory to GPU memory
cudaMemcpy(_data, _dataTocpy, _dataSize * sizeof(T),
cudaMemcpyDeviceToHost);
}
}
~moGPUCopy() {
}
};
#endif

74
trunk/paradiseo-gpu/src/memory/moGPUDeallocator.h Normal file
View file

@ -0,0 +1,74 @@
/*
<moGPUDeallocator.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Boufaras Karima, 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 __moGPUDeallocator_H_
#define __moGPUDeallocator_H_
/**
* class for Disallocation of data from GPU global memory
*/
class moGPUDeallocator {
public:
/**
* Constructor
*/
moGPUDeallocator() {
}
/**
*Deallocate data on GPU global memory
*@param _data the data to deallocate from GPU global memory
*/
template<typename T>
void operator()(T* & _data) {
//Deallocate data from GPU global memory
cudaFree(_data);
}
/**
* Destructor
*/
~moGPUDeallocator() {
}
};
#endif

96
trunk/paradiseo-gpu/src/memory/moGPUObject.h Normal file
View file

@ -0,0 +1,96 @@
/*
<moGPUObject.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Boufaras Karima, 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 __moGPUObject_H_
#define __moGPUObject_H_
#include <memory/moGPUAllocator.h>
#include <memory/moGPUDeallocator.h>
#include <memory/moGPUCopy.h>
/**
* class of data managment on GPU global memory (allocation,desallocation & copy)
*/
class moGPUObject {
public:
/*
* Constructor
*/
moGPUObject() {
}
/**
*Allocate & Copy data from CPU memory to GPU global memory (default copy)
*@param _data the data to allocate on GPU
*@param _dataTocpy the data to copy from CPU memory to _data on GPU memory
*@param _dataSize the size of data to copy
*/
template<typename T>
void memCopy(T* & _data, T * & _dataTocpy, unsigned _dataSize) {
malloc(_data, _dataSize);
copy(_data, _dataTocpy, _dataSize);
}
/**
*Copy device data from GPU global memory to global variable declared in device
*@param _dev_data the device global variable
*@param _dataTocpy the data to copy GPU global memory to GPU global variable
*/
template<typename T>
void memCopyGlobalVariable(T* & _dev_data, T * & _dataTocpy) {
copy(_dev_data, _dataTocpy);
}
/**
*Desallocate data on GPU global memory
*@param _data the data to desallocate from GPU global memory
*/
template<typename T>
void memFree(T* & _data) {
free(_data);
}
public:
moGPUAllocator malloc;
moGPUCopy copy;
moGPUDeallocator free;
};
#endif

85
trunk/paradiseo-gpu/src/memory/moGPUSpecificData.h Normal file
View file

@ -0,0 +1,85 @@
/*
<moGPUSpecificData.h>
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
Boufaras Karima, 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 __moGPUSpecificData_H_
#define __moGPUSpecificData_H_
#include <memory/moGPUObject.h>
/**
* class of managment of specific data problem
*/
class moGPUSpecificData {
public:
/*
* Default constructor
*/
moGPUSpecificData() {
sizeData = 0;
}
/*
* Load data from file given in argument
* @param _file the name of data file
*/
virtual void load(char* _file)=0;
/*
* Return the size
* @return the size of data to load
*/
virtual unsigned int getSize() {
return sizeData;
}
/*
* Set the size of data to load
*/
virtual void setSize(unsigned int _size) {
sizeData = _size;
}
public:
unsigned int sizeData;
moGPUObject GPUObject;
};
#endif

Some files were not shown because too many files have changed in this diff Show more