diff --git a/trunk/paradiseo-gpu/src/GPUType/moGPUBitVector.h b/trunk/paradiseo-gpu/src/GPUType/moGPUBitVector.h new file mode 100644 index 000000000..be50bc019 --- /dev/null +++ b/trunk/paradiseo-gpu/src/GPUType/moGPUBitVector.h @@ -0,0 +1,141 @@ +/* + + 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 +#include + +/** + * Implementation of Bit vector representation on GPU. + */ + +template + +class moGPUBitVector: public moGPUVector { + +public: + + /** + * Define bool vector corresponding to Solution + **/ + typedef bool ElemType; + using moGPUVector::vect; + using moGPUVector::N; + + /** + * Default constructor. + */ + + moGPUBitVector() : + moGPUVector () { + + } + + /** + *Constructor. + *@param _neighborhoodSize The neighborhood size. + */ + + moGPUBitVector(unsigned _neighborhoodSize) : + moGPUVector (_neighborhoodSize) { + create(); + } + + /** + *Constructor. + *@param _neighborhoodSize The neighborhood size. + *@param _b Value to assign to vector. + */ + + moGPUBitVector(unsigned _neighborhoodSize, bool _b) : + moGPUVector (_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 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 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::printOn(_os); + _os << ' '; + _os << N << ' '; + for (unsigned int i = 0; i < N; i++) + _os << (*this)[i] << ' '; + + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/GPUType/moGPUIntVector.h b/trunk/paradiseo-gpu/src/GPUType/moGPUIntVector.h new file mode 100644 index 000000000..a6b2a88ae --- /dev/null +++ b/trunk/paradiseo-gpu/src/GPUType/moGPUIntVector.h @@ -0,0 +1,122 @@ +/* + + 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 +#include + +/** + * Implementation of integer vector representation on GPU. + */ + +template + +class moGPUIntVector: public moGPUVector { + +public: + + using moGPUVector::vect; + using moGPUVector::N; + + /** + * Default constructor. + */ + + moGPUIntVector() : + moGPUVector () { + + } + + /** + *Constructor. + *@param _neighborhoodSize The neighborhood size. + */ + + moGPUIntVector(unsigned _neighborhoodSize) : + moGPUVector (_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 tmp_vect(_size); + if(_sizeN) { + for (unsigned i = 0; i ::printOn(_os); + _os << ' '; + _os << N << ' '; + unsigned int i; + for (i = 0; i < N; i++) + _os << vect[i] << ' '; + + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/GPUType/moGPUPermutationVector.h b/trunk/paradiseo-gpu/src/GPUType/moGPUPermutationVector.h new file mode 100644 index 000000000..bd3fcc4f3 --- /dev/null +++ b/trunk/paradiseo-gpu/src/GPUType/moGPUPermutationVector.h @@ -0,0 +1,92 @@ +/* + + 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 + +/** + * Implementation of Permutation vector representation on GPU. + */ + +template + +class moGPUPermutationVector: public moGPUIntVector { + +public: + + using moGPUIntVector::vect; + using moGPUIntVector::N; + + /** + * Default constructor. + */ + + moGPUPermutationVector() : + moGPUIntVector () { + + } + + /** + *Constructor. + *@param _size The solution size. + */ + + moGPUPermutationVector(unsigned _size) : + moGPUIntVector (_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 diff --git a/trunk/paradiseo-gpu/src/GPUType/moGPURealVector.h b/trunk/paradiseo-gpu/src/GPUType/moGPURealVector.h new file mode 100644 index 000000000..2a8b2bebe --- /dev/null +++ b/trunk/paradiseo-gpu/src/GPUType/moGPURealVector.h @@ -0,0 +1,123 @@ +/* + + 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 +#include + +/** + * Implementation of real vector representation on GPU. + */ + +template + +class moGPURealVector: public moGPUVector { + +public: + + using moGPUVector::vect; + using moGPUVector::N; + + /** + * Default constructor. + */ + + moGPURealVector() : + moGPUVector () { + + } + + /** + *Constructor. + *@param _neighborhoodSize The neighborhood size. + */ + + moGPURealVector(unsigned _neighborhoodSize) : + moGPUVector (_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 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 tmp_vect(_size); + for (unsigned i = 0; i ::printOn(_os); + _os << ' '; + _os << N << ' '; + unsigned int i; + for (i = 0; i < N; i++) + _os << vect[i] << ' '; + + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/GPUType/moGPUVector.h b/trunk/paradiseo-gpu/src/GPUType/moGPUVector.h new file mode 100644 index 000000000..d9b9afe83 --- /dev/null +++ b/trunk/paradiseo-gpu/src/GPUType/moGPUVector.h @@ -0,0 +1,184 @@ +/* + + 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 + +/** + * Implementation of a GPU solution representation. + */ + +template + +class moGPUVector: public EO { + +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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUEval.h b/trunk/paradiseo-gpu/src/eval/moGPUEval.h new file mode 100644 index 000000000..eca415860 --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUEval.h @@ -0,0 +1,146 @@ +/* + + 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 + +/** + * Abstract class for evaluation on GPU + */ + +template +class moGPUEval: public moEval { + +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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUEvalByCpy.h b/trunk/paradiseo-gpu/src/eval/moGPUEvalByCpy.h new file mode 100644 index 000000000..b15e32f44 --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUEvalByCpy.h @@ -0,0 +1,200 @@ +/* + + 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 +#include +#include + +/** + * class for the parallel evaluation of neighborhood + */ + +template +class moGPUEvalByCpy: public moGPUEval { + +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::neighborhoodSize; + using moGPUEval::host_FitnessArray; + using moGPUEval::device_FitnessArray; + using moGPUEval::device_solution; + using moGPUEval::NEW_kernel_Dim; + using moGPUEval::NEW_BLOCK_SIZE; + using moGPUEval::mutex; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval how to evaluate a neighbor + */ + + moGPUEvalByCpy(unsigned int _neighborhoodSize, Eval & _eval) : + moGPUEval (_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<<>>(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<<>>(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<<>>(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<<>>(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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUEvalByModif.h b/trunk/paradiseo-gpu/src/eval/moGPUEvalByModif.h new file mode 100644 index 000000000..935faa90a --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUEvalByModif.h @@ -0,0 +1,199 @@ +/* + + 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 +#include +#include + +/** + * class for the parallel evaluation of neighborhood + */ + +template +class moGPUEvalByModif: public moGPUEval { + +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::neighborhoodSize; + using moGPUEval::host_FitnessArray; + using moGPUEval::device_FitnessArray; + using moGPUEval::device_solution; + using moGPUEval::NEW_kernel_Dim; + using moGPUEval::NEW_BLOCK_SIZE; + using moGPUEval::mutex; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval the incremental evaluation + */ + + moGPUEvalByModif(unsigned int _neighborhoodSize, Eval & _eval) : + moGPUEval (_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<<>>(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<<>>(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<<>>(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<<>>(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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUEvalFunc.h b/trunk/paradiseo-gpu/src/eval/moGPUEvalFunc.h new file mode 100644 index 000000000..25eb907fa --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUEvalFunc.h @@ -0,0 +1,84 @@ +/* + + 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 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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUKernelEvalByCpy.h b/trunk/paradiseo-gpu/src/eval/moGPUKernelEvalByCpy.h new file mode 100644 index 000000000..858b132ac --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUKernelEvalByCpy.h @@ -0,0 +1,71 @@ +/* + + 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 + +__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 + 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 + +__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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByCpy.h b/trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByCpy.h new file mode 100644 index 000000000..46a883f25 --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByCpy.h @@ -0,0 +1,202 @@ +/* + + 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 +#include +#include + +/** + * class for the Mapping neighborhood evaluation + */ + +template +class moGPUMappingEvalByCpy: public moGPUEval { + +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::neighborhoodSize; + using moGPUEval::host_FitnessArray; + using moGPUEval::device_FitnessArray; + using moGPUEval::device_solution; + using moGPUEval::NEW_BLOCK_SIZE; + using moGPUEval::NEW_kernel_Dim; + using moGPUEval::mutex; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval how to evaluate a neighbor + */ + + moGPUMappingEvalByCpy(unsigned int _neighborhoodSize, Eval & _eval) : + moGPUEval (_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<<>>(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<<>>(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<<>>(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<<>>(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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByModif.h b/trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByModif.h new file mode 100644 index 000000000..eba3a5fa4 --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUMappingEvalByModif.h @@ -0,0 +1,200 @@ +/* + + 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 +#include +#include + +/** + * class for the Mapping neighborhood evaluation + */ + +template +class moGPUMappingEvalByModif: public moGPUEval { + +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::neighborhoodSize; + using moGPUEval::host_FitnessArray; + using moGPUEval::device_FitnessArray; + using moGPUEval::device_solution; + using moGPUEval::NEW_BLOCK_SIZE; + using moGPUEval::NEW_kernel_Dim; + using moGPUEval::mutex; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval the incremental evaluation + */ + + moGPUMappingEvalByModif(unsigned int _neighborhoodSize, Eval & _eval) : + moGPUEval (_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<<>>(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<<>>(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<<>>(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<<>>(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 diff --git a/trunk/paradiseo-gpu/src/eval/moGPUMappingKernelEvalByCpy.h b/trunk/paradiseo-gpu/src/eval/moGPUMappingKernelEvalByCpy.h new file mode 100644 index 000000000..7d204c337 --- /dev/null +++ b/trunk/paradiseo-gpu/src/eval/moGPUMappingKernelEvalByCpy.h @@ -0,0 +1,76 @@ +/* + + 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 + +__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 + 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 + +__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 + 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 + void operator()(T* & _data, unsigned _dataSize) { + + //Allocate data in GPU memory + cudaMalloc((void**) &_data, _dataSize * sizeof(T)); + + } + + /** + * Destructor + */ + + ~moGPUAllocator() { + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/memory/moGPUCopy.h b/trunk/paradiseo-gpu/src/memory/moGPUCopy.h new file mode 100644 index 000000000..2d7c4b284 --- /dev/null +++ b/trunk/paradiseo-gpu/src/memory/moGPUCopy.h @@ -0,0 +1,112 @@ +/* + + 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 + 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 + 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 + 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 diff --git a/trunk/paradiseo-gpu/src/memory/moGPUDeallocator.h b/trunk/paradiseo-gpu/src/memory/moGPUDeallocator.h new file mode 100644 index 000000000..18ef65d9c --- /dev/null +++ b/trunk/paradiseo-gpu/src/memory/moGPUDeallocator.h @@ -0,0 +1,74 @@ +/* + + 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 + void operator()(T* & _data) { + + //Deallocate data from GPU global memory + cudaFree(_data); + + } + + /** + * Destructor + */ + + ~moGPUDeallocator() { + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/memory/moGPUObject.h b/trunk/paradiseo-gpu/src/memory/moGPUObject.h new file mode 100644 index 000000000..74f2ccb7f --- /dev/null +++ b/trunk/paradiseo-gpu/src/memory/moGPUObject.h @@ -0,0 +1,96 @@ +/* + + 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 +#include +#include + +/** + * 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 + 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 + 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 + void memFree(T* & _data) { + free(_data); + } + +public: + + moGPUAllocator malloc; + moGPUCopy copy; + moGPUDeallocator free; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/memory/moGPUSpecificData.h b/trunk/paradiseo-gpu/src/memory/moGPUSpecificData.h new file mode 100644 index 000000000..b3286dadc --- /dev/null +++ b/trunk/paradiseo-gpu/src/memory/moGPUSpecificData.h @@ -0,0 +1,85 @@ +/* + + 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 + +/** + * 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 diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUBitNeighbor.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUBitNeighbor.h new file mode 100644 index 000000000..b48753bac --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUBitNeighbor.h @@ -0,0 +1,122 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, 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 __moGPUBitNeighbor_h +#define __moGPUBitNeighbor_h + +#include +#include +#include + +/** + * Neighbor related to a solution vector of Bit + */ + +template +class moGPUBitNeighbor: public moBackableNeighbor > , + public moIndexNeighbor > { + +public: + + typedef moGPUBitVector EOT ; + using moBackableNeighbor::fitness; + using moIndexNeighbor::key; + + /** + * move the solution + * @param _solution the solution to move + */ + + virtual void move(EOT & _solution) { + + _solution[key] = !_solution[key]; + _solution.invalidate(); + + } + + /** + * move back the solution (useful for the evaluation by modif) + * @param _solution the solution to move back + */ + + virtual void moveBack(EOT & _solution) { + + move(_solution); + + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUBitNeighbor"; + } + + /** + * Read object.\ + * Calls base class, just in case that one had something to do. + * The read and print methods should be compatible and have the same format. + * In principle, format is "plain": they just print a number + * @param _is a std::istream. + * @throw runtime_std::exception If a valid object can't be read. + */ + + virtual void readFrom(std::istream& _is) { + std::string fitness_str; + int pos = _is.tellg(); + _is >> fitness_str; + if (fitness_str == "INVALID") { + throw std::runtime_error("invalid fitness"); + } else { + Fitness repFit; + _is.seekg(pos); + _is >> repFit; + _is >> key; + fitness(repFit); + } + } + + /** + * Write object. Called printOn since it prints the object _on_ a stream. + * @param _os A std::ostream. + */ + + virtual void printOn(std::ostream& _os) const { + _os << fitness() << ' ' << key << std::endl; + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhood.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhood.h new file mode 100644 index 000000000..08a94d31b --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhood.h @@ -0,0 +1,110 @@ +/* + + 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 ue, + 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". + + 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 __moGPUMappingNeighborhood_h +#define __moGPUMappingNeighborhood_h + +#include +#include + +template +class moGPUMappingNeighborhood: public moMappingNeighborhood { + +public: + + /** + * Define a Neighbor and type of a solution corresponding + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + using moMappingNeighborhood::neighborhoodSize; + using moMappingNeighborhood::currentIndex; + using moMappingNeighborhood::indices; + using moMappingNeighborhood::mapping; + using moMappingNeighborhood::xChange; + using moMappingNeighborhood::mutex; + + /** + * Constructor + * @param _neighborhoodSize the neighborhood size + * @param _xChange the number of x-change positions + */ + + moGPUMappingNeighborhood(unsigned int _neighborhoodSize, + unsigned int _xChange) : + moMappingNeighborhood (_neighborhoodSize, _xChange) { + sendMapping = false; + cudaMalloc((void**) &device_Mapping, sizeof(unsigned int) + * neighborhoodSize * _xChange); + } + ; + + /** + *Destructor + */ + + ~moGPUMappingNeighborhood() { + + cudaFree(device_Mapping); + } + + /** + * Initialization of the neighborhood and mapping on device + * @param _solution the solution to explore + * @param _current the first neighbor + */ + + virtual void init(EOT& _solution, Neighbor& _current) { + + moMappingNeighborhood::init(_solution, _current); + if (!sendMapping) { + cudaMemcpy(device_Mapping, mapping,xChange * neighborhoodSize + * sizeof(unsigned int), cudaMemcpyHostToDevice); + sendMapping = true; + } + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUMappingNeighborhood"; + } + +protected: + + bool sendMapping; + unsigned int * device_Mapping; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhoodByCpy.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhoodByCpy.h new file mode 100644 index 000000000..68031f6bf --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhoodByCpy.h @@ -0,0 +1,97 @@ +/* + + 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 ue, + 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". + + 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 __moGPUMappingNeighborhoodByCpy_h +#define __moGPUMappingNeighborhoodByCpy_h + +#include +#include + +template +class moGPUMappingNeighborhoodByCpy: public moGPUMappingNeighborhood { + +public: + + /** + * Define a Neighbor and type of a solution corresponding + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + /*A tester*/ + using moGPUMappingNeighborhood::neighborhoodSize; + using moGPUMappingNeighborhood::currentIndex; + using moGPUMappingNeighborhood::indices; + using moGPUMappingNeighborhood::mapping; + using moGPUMappingNeighborhood::xChange; + using moGPUMappingNeighborhood::mutex; + using moGPUMappingNeighborhood::device_Mapping; + + /** + * Constructor + * @param _neighborhoodSize the neighborhood size + * @param _xChange the number of x-change positions + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + + moGPUMappingNeighborhoodByCpy(unsigned int _neighborhoodSize, + unsigned int _xChange,moGPUEval& _eval) : + moGPUMappingNeighborhood (_neighborhoodSize, _xChange), eval(_eval){ + } + + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _current the first neighbor + */ + + virtual void init(EOT& _solution, Neighbor& _current) { + + moGPUMappingNeighborhood::init(_solution, _current); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution, device_Mapping,1,1); + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUMappingNeighborhoodByCpy"; + } + +protected: + + moGPUEval& eval; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhoodByModif.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhoodByModif.h new file mode 100644 index 000000000..c1bab360e --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUMappingNeighborhoodByModif.h @@ -0,0 +1,97 @@ +/* + + 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 ue, + 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". + + 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 __moGPUMappingNeighborhoodByModif_h +#define __moGPUMappingNeighborhoodByModif_h + +#include +#include + +template +class moGPUMappingNeighborhoodByModif: public moGPUMappingNeighborhood { + +public: + + /** + * Define a Neighbor and type of a solution corresponding + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + /*A tester*/ + using moGPUMappingNeighborhood::neighborhoodSize; + using moGPUMappingNeighborhood::currentIndex; + using moGPUMappingNeighborhood::indices; + using moGPUMappingNeighborhood::mapping; + using moGPUMappingNeighborhood::xChange; + using moGPUMappingNeighborhood::mutex; + using moGPUMappingNeighborhood::device_Mapping; + + /** + * Constructor + * @param _neighborhoodSize the neighborhood size + * @param _xChange the number of x-change positions + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + + moGPUMappingNeighborhoodByModif(unsigned int _neighborhoodSize, + unsigned int _xChange,moGPUEval& _eval) : + moGPUMappingNeighborhood (_neighborhoodSize, _xChange), eval(_eval){ + } + + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _current the first neighbor + */ + + virtual void init(EOT& _solution, Neighbor& _current) { + + moGPUMappingNeighborhood::init(_solution, _current); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution, device_Mapping,0,1); + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUMappingNeighborhoodByModif"; + } + +protected: + + moGPUEval& eval; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUOrderNeighborhoodByCpy.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUOrderNeighborhoodByCpy.h new file mode 100644 index 000000000..bce7abfb0 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUOrderNeighborhoodByCpy.h @@ -0,0 +1,96 @@ +/* + + 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 __moGPUOrderNeighborhoodByCpy_h +#define __moGPUOrderNeighborhoodByCpy_h + +#include +#include + +/** + * An ordered neighborhood with parallel evaluation + */ + +template +class moGPUOrderNeighborhoodByCpy: public moOrderNeighborhood { + + public: + + /** + * Define type of a solution corresponding to Neighbor + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + using moOrderNeighborhood::neighborhoodSize; + using moOrderNeighborhood::currentIndex; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + + moGPUOrderNeighborhoodByCpy(unsigned int _neighborhoodSize, + moGPUEval& _eval) : + moOrderNeighborhood (_neighborhoodSize), eval(_eval) { + } + + /** + * Initialization of the neighborhood + *@param _solution the solution to explore + *@param _neighbor the first neighbor + */ + + virtual void init(EOT & _solution, Neighbor & _neighbor) { + + moOrderNeighborhood::init(_solution, _neighbor); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution,1,1); + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUOrderNeighborhoodByCpy"; + } + + protected: + moGPUEval& eval; +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUOrderNeighborhoodByModif.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUOrderNeighborhoodByModif.h new file mode 100644 index 000000000..7c997a901 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUOrderNeighborhoodByModif.h @@ -0,0 +1,96 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, Thé Van LUONG, Karima Boufaras + + This software is governed by the CeCILL license under French law and + abiding by the rules of distribution of free software. You can use, + modify and/ or redistribute the software under the terms of the CeCILL + license as circulated by CEA, CNRS and INRIA at the following URL + "http://www.cecill.info". + + As a counterpart to the access to the source code and rights to copy, + modify and redistribute granted by the license, users are provided only + with a limited warranty and the software's author, the holder of the + economic rights, and the successive licensors have only limited liability. + + In this respect, the user's attention is drawn to the risks associated + with loading, using, modifying and/or developing or reproducing the + software by the user in light of its specific status of free software, + that may mean that it is complicated to manipulate, and that also + therefore means that it is reserved for developers and experienced + professionals having in-depth computer knowledge. Users are therefore + encouraged to load and test the software's suitability as regards their + requirements in conditions enabling the security of their systems and/or + data to be ensured and, more generally, to use and operate it in the + same conditions as regards security. + The fact that you are presently reading this means that you have had + knowledge of the CeCILL license and that you accept its terms. + + ParadisEO WebSite : http://paradiseo.gforge.inria.fr + Contact: paradiseo-help@lists.gforge.inria.fr +*/ + +#ifndef __moGPUOrderNeighborhoodByModif_h +#define __moGPUOrderNeighborhoodByModif_h + +#include +#include + +/** + * An ordered neighborhood with parallel evaluation + */ + +template +class moGPUOrderNeighborhoodByModif: public moOrderNeighborhood { + + public: + + /** + * Define type of a solution corresponding to Neighbor + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + using moOrderNeighborhood::neighborhoodSize; + using moOrderNeighborhood::currentIndex; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + + moGPUOrderNeighborhoodByModif(unsigned int _neighborhoodSize, + moGPUEval& _eval) : + moOrderNeighborhood (_neighborhoodSize), eval(_eval) { + } + + /** + * Initialization of the neighborhood + *@param _solution the solution to explore + *@param _neighbor the first neighbor + */ + + virtual void init(EOT & _solution, Neighbor & _neighbor) { + + moOrderNeighborhood::init(_solution, _neighbor); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution,0,1); + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUOrderNeighborhoodByModif"; + } + + protected: + moGPUEval& eval; +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithReplNeighborhoodByCpy.h b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithReplNeighborhoodByCpy.h new file mode 100644 index 000000000..8c99c82b3 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithReplNeighborhoodByCpy.h @@ -0,0 +1,88 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, 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 _moGPURndWithReplNeighborhood_h +#define _moGPURndWithReplNeighborhood_h + +#include +#include + +/** + * A Random With replacement Neighborhood with parallel evaluation + */ +template +class moGPURndWithReplNeighborhoodByCpy: public moRndWithReplNeighborhood { + public: + + /** + * Define type of a solution corresponding to Neighbor + */ + typedef typename Neighbor::EOT EOT; + + using moRndWithReplNeighborhood::neighborhoodSize; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + moGPURndWithReplNeighborhoodByCpy(unsigned int _neighborhoodSize, moGPUEvalByCpy< + Neighbor>& _eval) : + moRndWithReplNeighborhood (_neighborhoodSize), eval(_eval) { + } + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _neighbor the first neighbor + */ + virtual void init(EOT & _solution, Neighbor & _neighbor) { + moRndWithReplNeighborhood::init(_solution, _neighbor); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution,0,1); + + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPURndWithReplNeighborhoodByCpy"; + } + protected: + moGPUEvalByCpy& eval; +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithReplNeighborhoodByModif.h b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithReplNeighborhoodByModif.h new file mode 100644 index 000000000..fc64d3288 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithReplNeighborhoodByModif.h @@ -0,0 +1,87 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, 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 __moGPURndWithReplNeighborhoodByModif_h +#define __moGPURndWithReplNeighborhoodByModif_h + +#include +#include + +/** + * A Random With replacement Neighborhood with parallel evaluation + */ +template +class moGPURndWithReplNeighborhoodByModif: public moRndWithReplNeighborhood { + public: + + /** + * Define type of a solution corresponding to Neighbor + */ + typedef typename Neighbor::EOT EOT; + + using moRndWithReplNeighborhood::neighborhoodSize; + + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + moGPURndWithReplNeighborhoodByModif(unsigned int _neighborhoodSize, moGPUEval& _eval) : + moRndWithReplNeighborhood (_neighborhoodSize), eval(_eval) { + } + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _neighbor the first neighbor + */ + virtual void init(EOT & _solution, Neighbor & _neighbor) { + moRndWithReplNeighborhood::init(_solution, _neighbor); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution,0,1); + + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPURndWithReplNeighborhood"; + } + protected: + moGPUEval& eval; +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithoutReplNeighborhoodByCpy.h b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithoutReplNeighborhoodByCpy.h new file mode 100644 index 000000000..69108b8d2 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithoutReplNeighborhoodByCpy.h @@ -0,0 +1,90 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, 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 __moGPURndWithoutReplNeighborhood_h +#define __moGPURndWithoutReplNeighborhood_h + +#include +#include + +/** + * A Random without replacement Neighborhood with parallel evaluation + */ +template +class moGPURndWithoutReplNeighborhood: public moRndWithoutReplNeighborhood { + public: + + /** + * Define type of a solution corresponding to Neighbor + */ + typedef typename Neighbor::EOT EOT; + + using moRndWithoutReplNeighborhood::neighborhoodSize; + using moRndWithoutReplNeighborhood::maxIndex; + using moRndWithoutReplNeighborhood::indexVector; + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + moGPURndWithoutReplNeighborhood(unsigned int _neighborhoodSize,moGPUEval< + Neighbor>& _eval) : + moRndWithoutReplNeighborhood (_neighborhoodSize),eval(_eval) { + for (unsigned int i = 0; i < neighborhoodSize; i++) + indexVector.push_back(i); + } + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _neighbor the first neighbor + */ + virtual void init(EOT & _solution, Neighbor & _neighbor) { + moRndWithoutReplNeighborhood::init(_solution, _neighbor); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution,0,1); + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPURndWithoutReplNeighborhood"; + } + + protected: + moGPUEval& eval; +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithoutReplNeighborhoodByModif.h b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithoutReplNeighborhoodByModif.h new file mode 100644 index 000000000..f1b4ff1cb --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPURndWithoutReplNeighborhoodByModif.h @@ -0,0 +1,90 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, 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 __moGPURndWithoutReplNeighborhood_h +#define __moGPURndWithoutReplNeighborhood_h + +#include +#include + +/** + * A Random without replacement Neighborhood with parallel evaluation + */ +template +class moGPURndWithoutReplNeighborhood: public moRndWithoutReplNeighborhood { + public: + + /** + * Define type of a solution corresponding to Neighbor + */ + typedef typename Neighbor::EOT EOT; + + using moRndWithoutReplNeighborhood::neighborhoodSize; + using moRndWithoutReplNeighborhood::maxIndex; + using moRndWithoutReplNeighborhood::indexVector; + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + moGPURndWithoutReplNeighborhood(unsigned int _neighborhoodSize,moGPUEvalByModif< + Neighbor>& _eval) : + moRndWithoutReplNeighborhood (_neighborhoodSize),eval(_eval) { + for (unsigned int i = 0; i < neighborhoodSize; i++) + indexVector.push_back(i); + } + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _neighbor the first neighbor + */ + virtual void init(EOT & _solution, Neighbor & _neighbor) { + moRndWithoutReplNeighborhood::init(_solution, _neighbor); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution,0,1); + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPURndWithoutReplNeighborhood"; + } + + protected: + moGPUEvalByModif& eval; +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUXBitFlippingNeighbor.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUXBitFlippingNeighbor.h new file mode 100644 index 000000000..4cb649a65 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUXBitFlippingNeighbor.h @@ -0,0 +1,101 @@ +/* + + 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 ue, + 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". + + 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 __moGPUXBitFlippingNeighbor_h +#define __moGPUXBitFlippingNeighbor_h + +#include +#include +#include + +/** + * A GPU X-BitFlipping Neighbor + */ + +template< class Fitness > +class moGPUXBitFlippingNeighbor:public moBackableNeighbor< moGPUBitVector > , + public moXChangeNeighbor< moGPUBitVector > { +public: + + typedef moGPUBitVector EOT ; + using moXChangeNeighbor::indices; + using moXChangeNeighbor::xChange; + using moXChangeNeighbor::key; + + /** + *Default Constructor + */ + + moGPUXBitFlippingNeighbor() : + moXChangeNeighbor () { + } + + /** + * Constructor + * @param _xFlip the number of bit to flip + */ + + moGPUXBitFlippingNeighbor(unsigned int _xFlip) : + moXChangeNeighbor (_xFlip) { + } + + /** + * Apply the K-Flip in solution + * @param _solution the solution to move + */ + + virtual void move(EOT& _solution) { + for (unsigned int i = 0; i < xChange; i++) + _solution[indices[i]] = !_solution[indices[i]]; + _solution.invalidate(); + + } + + /** + * apply the K-Flip to restore the solution (use by moFullEvalByModif) + * @param _solution the solution to move back + */ + + virtual void moveBack(EOT& _solution) { + move(_solution); + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + + virtual std::string className() const { + return "moGPUXBitFlippingNeighbor"; + } + +}; + +#endif + diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhood.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhood.h new file mode 100644 index 000000000..1ea85aee7 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhood.h @@ -0,0 +1,110 @@ +/* + + 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 ue, + 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". + + 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 __moGPUXChangeNeighborhood_h +#define __moGPUXChangeNeighborhood_h + +#include +#include + +template +class moGPUXChangeNeighborhood: public moXChangeNeighborhood { + +public: + + /** + * Define a Neighbor and type of a solution corresponding + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + using moXChangeNeighborhood::neighborhoodSize; + using moXChangeNeighborhood::currentIndex; + using moXChangeNeighborhood::indices; + using moXChangeNeighborhood::mapping; + using moXChangeNeighborhood::xChange; + using moXChangeNeighborhood::mutex; + + /** + * Constructor + * @param _neighborhoodSize the neighborhood size + * @param _xChange the number of x-change positions + */ + + moGPUXChangeNeighborhood(unsigned int _neighborhoodSize, + unsigned int _xChange) : + moXChangeNeighborhood (_neighborhoodSize, _xChange) { + sendMapping = false; + cudaMalloc((void**) &device_Mapping, sizeof(unsigned int) + * neighborhoodSize * _xChange); + } + + + /** + *Destructor + */ + + ~moGPUXChangeNeighborhood() { + + cudaFree(device_Mapping); + } + + /** + * Initialization of the neighborhood and mapping on device + * @param _solution the solution to explore + * @param _current the first neighbor + */ + + virtual void init(EOT& _solution, Neighbor& _current) { + + moXChangeNeighborhood::init(_solution, _current); + if (!sendMapping) { + cudaMemcpy(device_Mapping, mapping, xChange * neighborhoodSize + * sizeof(unsigned int), cudaMemcpyHostToDevice); + sendMapping = true; + } + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUXChangeNeighborhood"; + } + +protected: + + bool sendMapping; + unsigned int * device_Mapping; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhoodByCpy.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhoodByCpy.h new file mode 100644 index 000000000..b417144b3 --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhoodByCpy.h @@ -0,0 +1,97 @@ +/* + + 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 ue, + 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". + + 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 __moGPUXChangeNeighborhoodByCpy_h +#define __moGPUXChangeNeighborhoodByCpy_h + +#include +#include + +template +class moGPUXChangeNeighborhoodByCpy: public moGPUXChangeNeighborhood { + +public: + + /** + * Define a Neighbor and type of a solution corresponding + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + /*A tester*/ + using moGPUXChangeNeighborhood::neighborhoodSize; + using moGPUXChangeNeighborhood::currentIndex; + using moGPUXChangeNeighborhood::indices; + using moGPUXChangeNeighborhood::mapping; + using moGPUXChangeNeighborhood::xChange; + using moGPUXChangeNeighborhood::mutex; + using moGPUXChangeNeighborhood::device_Mapping; + + /** + * Constructor + * @param _neighborhoodSize the neighborhood size + * @param _xChange the number of x-change positions + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + + moGPUXChangeNeighborhoodByCpy(unsigned int _neighborhoodSize, + unsigned int _xChange,moGPUEval& _eval) : + moGPUXChangeNeighborhood (_neighborhoodSize, _xChange), eval(_eval){ + } + + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _current the first neighbor + */ + + virtual void init(EOT& _solution, Neighbor& _current) { + + moGPUXChangeNeighborhood::init(_solution, _current); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution, device_Mapping,1,1); + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUXChangeNeighborhoodByCpy"; + } + +protected: + + moGPUEval& eval; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhoodByModif.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhoodByModif.h new file mode 100644 index 000000000..c1cc224ad --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUXChangeNeighborhoodByModif.h @@ -0,0 +1,96 @@ +/* + + 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 ue, + 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". + + 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 __moGPUXChangeNeighborhoodByModif_h +#define __moGPUXChangeNeighborhoodByModif_h + +#include +#include + +template +class moGPUXChangeNeighborhoodByModif: public moGPUXChangeNeighborhood { + +public: + + /** + * Define a Neighbor and type of a solution corresponding + */ + + typedef N Neighbor; + typedef typename Neighbor::EOT EOT; + + using moGPUXChangeNeighborhood::neighborhoodSize; + using moGPUXChangeNeighborhood::currentIndex; + using moGPUXChangeNeighborhood::indices; + using moGPUXChangeNeighborhood::mapping; + using moGPUXChangeNeighborhood::xChange; + using moGPUXChangeNeighborhood::mutex; + using moGPUXChangeNeighborhood::device_Mapping; + + /** + * Constructor + * @param _neighborhoodSize the neighborhood size + * @param _xChange the number of x-change positions + * @param _eval show how to evaluate neighborhood of a solution at one time + */ + + moGPUXChangeNeighborhoodByModif(unsigned int _neighborhoodSize, + unsigned int _xChange,moGPUEval& _eval) : + moGPUXChangeNeighborhood (_neighborhoodSize, _xChange), eval(_eval){ + } + + + /** + * Initialization of the neighborhood + * @param _solution the solution to explore + * @param _current the first neighbor + */ + + virtual void init(EOT& _solution, Neighbor& _current) { + + moGPUXChangeNeighborhood::init(_solution, _current); + //Compute all neighbors fitness at one time + eval.neighborhoodEval(_solution, device_Mapping,0,1); + } + + /** + * Return the class Name + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUXChangeNeighborhoodByModif"; + } + +protected: + + moGPUEval& eval; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/neighborhood/moGPUXSwapNeighbor.h b/trunk/paradiseo-gpu/src/neighborhood/moGPUXSwapNeighbor.h new file mode 100644 index 000000000..5ddb2cb2a --- /dev/null +++ b/trunk/paradiseo-gpu/src/neighborhood/moGPUXSwapNeighbor.h @@ -0,0 +1,108 @@ +/* + + 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 ue, + 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". + + 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 __moGPUXSwapNeighbor_h +#define __moGPUXSwapNeighbor_h +#include +#include +#include + +/** + * A GPU X-Swap Neighbor + */ + +template +class moGPUXSwapNeighbor: public moBackableNeighbor > , + public moXChangeNeighbor > { + +public: + + typedef moGPUPermutationVector EOT ; + using moXChangeNeighbor::indices; + using moXChangeNeighbor::xChange; + using moXChangeNeighbor::key; + + /** + *Default Constructor + */ + + moGPUXSwapNeighbor() : + moXChangeNeighbor () { + } + + /** + * Constructor + * @param _xSwap the number of swap to do + */ + + moGPUXSwapNeighbor(unsigned int _xSwap) : + moXChangeNeighbor (_xSwap) { + } + + /** + * Apply the K-swap + * @param _solution the solution to move + */ + virtual void move(EOT& _solution) { + EOT tmp(1); + for (unsigned int i = 0; i < xChange-1; i++) { + tmp[0] = _solution[indices[i]]; + _solution[indices[i]] = _solution[indices[i + 1]]; + _solution[indices[i + 1]] = tmp[0]; + } + _solution.invalidate(); + + } + + /** + * apply the K-swap to restore the solution (use by moFullEvalByModif) + * @param _solution the solution to move back + */ + virtual void moveBack(EOT& _solution) { + EOT tmp(1); + for (int i = xChange-1; i > 0; i--) { + tmp[0] = _solution[indices[i]]; + _solution[indices[i]] = _solution[indices[i - 1]]; + _solution[indices[i - 1]] = tmp[0]; + } + _solution.invalidate(); + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUXSwapNeighbor"; + } + +}; + +#endif + diff --git a/trunk/paradiseo-gpu/src/performance/moGPUTimer.h b/trunk/paradiseo-gpu/src/performance/moGPUTimer.h new file mode 100644 index 000000000..44b3dd10c --- /dev/null +++ b/trunk/paradiseo-gpu/src/performance/moGPUTimer.h @@ -0,0 +1,108 @@ +/* + + 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 __moGPUTimer_H_ +#define __moGPUTimer_H_ + +#include + +/** + * To compute execution time + */ + +class moGPUTimer { + +public: + + timeval tim; + double t1; + double t2; + + /** + * Constructor + */ + + moGPUTimer() { + + t1 = 0; + t2 = 0; + } + + /** + * Destructor + */ + + ~moGPUTimer() { + + } + + /** + * to start compute execution time + */ + + void start() { + + // Start timer + gettimeofday(&tim, NULL); + t1 = tim.tv_sec + (tim.tv_usec / 1000000.0); + } + + /** + * to stop compute execution time + */ + + void stop() { + + // Stop timer + gettimeofday(&tim, NULL); + t2 = tim.tv_sec + (tim.tv_usec / 1000000.0); + + } + + /** + * to get timer value in s + * @return execution time in s + */ + + double getTime() { + + // get time + return (t2 - t1); + + + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/problems/data/PPPData.h b/trunk/paradiseo-gpu/src/problems/data/PPPData.h new file mode 100644 index 000000000..3be2c103e --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/data/PPPData.h @@ -0,0 +1,171 @@ +/* + + 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 _PPPData_H_ +#define _PPPData_H_ + +#include + +template +class PPPData: public moGPUSpecificData { + +public: + + using moGPUSpecificData::GPUObject; + + /** + * Default Constructor + */ + + PPPData() : + moGPUSpecificData() { + + //(*this).load(); + } + + /** + * Constructor by copy + * @param _pppData the specific data of PPP + */ + + PPPData(const PPPData & _pppData) { + + a_h = new int[Md * Nd]; + H_h = new int[Nd]; + + for (int i = 0; i < Md; i++) + for (int j = 0; j < Nd; j++) { + a_h[i * Nd + j] = _pppData.a_h[i * Nd + j]; + } + for (int k = 0; k < Nd; k++) { + H_h[k] = _pppData.H_h[k]; + } + + GPUObject.memCopy(a_d, a_h, Nd * Md); + GPUObject.memCopy(H_d, H_h, Nd); + + } + + /** + * Assignement operator + * @param _pppData the specific data of PPP + * @return a PPP Data + */ + + PPPData & operator=(const PPPData & _pppData) { + + a_h = new int[Md * Nd]; + H_h = new int[Nd]; + for (int i = 0; i < Md; i++) + for (int j = 0; j < Nd; j++) { + a_h[i * Nd + j] = _pppData.a_h[i * Nd + j]; + } + for (int k = 0; k < Nd; k++) { + H_h[k] = _pppData.H_h[k]; + } + + GPUObject.memCopy(a_d, a_h, Nd * Md); + GPUObject.memCopy(H_d, H_h, Nd); + + return (*this); + } + + /* + * Destructor + */ + + ~PPPData() { + GPUObject.memFree(a_d); + GPUObject.memFree(H_d); + delete[] a_h; + delete[] H_h; + } + + /* + *Load PPP data + */ + + void load(char * _fileName) { + } + + void load() { + + int *v = new int[Nd]; + int *s = new int[Md]; + a_h = new int[Md * Nd]; + H_h = new int[Nd]; + for (int i = 0; i < Nd; i++) + H_h[i] = 0; + for (int i = 0; i < Md; i++) { + for (int j = 0; j < Nd; j++) { + a_h[i * Nd + j] =pow(-1,rand()); + } + } + for (int i = 0; i < Nd; i++) { + v[i]=pow(-1,rand()); + } + + for (int i = 0; i < Md; i++) { + s[i] = 0; + for (int j = 0; j < Nd; j++) + s[i] += a_h[i * Nd + j] * v[j]; + if (s[i] < 0) { + for (int k = 0; k < Nd; k++) + a_h[i * Nd + k] = -a_h[i * Nd + k]; + s[i] = -s[i]; + } + if(s[i]>0) + H_h[s[i]-1]++; + } + + //Allocate and copy QAP data from CPU memory to GPU global memory + GPUObject.memCopy(a_d, a_h, Nd * Md); + GPUObject.memCopy(H_d, H_h, Nd); + + delete[] v; + delete[] s; + + } + + +public: + + ElemType* a_h; + ElemType* H_h; + ElemType* a_d; + ElemType* H_d; + +}; +#endif + diff --git a/trunk/paradiseo-gpu/src/problems/data/QAPData.h b/trunk/paradiseo-gpu/src/problems/data/QAPData.h new file mode 100644 index 000000000..60d079c7a --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/data/QAPData.h @@ -0,0 +1,167 @@ +/* + + 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 _QAPData_H_ +#define _QAPData_H_ + +#include + +template +class QAPData: public moGPUSpecificData { + +public: + + using moGPUSpecificData::sizeData; + using moGPUSpecificData::GPUObject; + + /** + * Default Constructor + */ + + QAPData() : + moGPUSpecificData() { + } + + /** + * Constructor + * @param _fileName the data file name + */ + + QAPData(char* _fileName) { + + (*this).load(_fileName); + + } + + /** + * Constructor by copy + * @param _qapData the specific data of QAP + */ + + QAPData(const QAPData & _qapData) { + + sizeData = _qapData.sizeData; + a_h = new int[sizeData * sizeData]; + b_h = new int[sizeData * sizeData]; + for (int i = 0; i < sizeData; i++) + + for (int j = 0; j < sizeData; j++) { + + a_h[i * sizeData + j] = _qapData.a_h[i * sizeData + j]; + b_h[i * sizeData + j] = _qapData.b_h[i * sizeData + j]; + + } + + GPUObject.memCopy(a_d, a_h, sizeData * sizeData); + GPUObject.memCopy(b_d, b_h, sizeData * sizeData); + + } + + /** + * Assignement operator + * @param _qapData the specific data of QAP + * @return a QAP Data + */ + + QAPData & operator=(const QAPData & _qapData) { + + sizeData = _qapData.sizeData; + a_h = new int[sizeData * sizeData]; + b_h = new int[sizeData * sizeData]; + + for (int i = 0; i < sizeData; i++) + for (int j = 0; j < sizeData; j++) { + + a_h[i * sizeData + j] = _qapData.a_h[i * sizeData + j]; + b_h[i * sizeData + j] = _qapData.b_h[i * sizeData + j]; + + } + GPUObject.memCopy(a_d, a_h, sizeData * sizeData); + GPUObject.memCopy(b_d, b_h, sizeData * sizeData); + return (*this); + } + + /* + * Destructor + */ + + ~QAPData() { + GPUObject.memFree(a_d); + GPUObject.memFree(b_d); + delete[] a_h; + delete[] b_h; + } + + /* + *Load QAP data from file name + *@param _fileName the data file name to load + */ + + void load(char* _fileName) { + + fstream file(_fileName, ios::in); + if (!file) { + + string str = "QAPData: Could not open file [" + (string) _fileName + + "]."; + throw runtime_error(str); + } + + unsigned i, j; + file >> sizeData; + a_h = new ElemType[sizeData * sizeData]; + b_h = new ElemType[sizeData * sizeData]; + + for (i = 0; i < sizeData; i++) + for (j = 0; j < sizeData; j++) + file >> a_h[i * sizeData + j]; + for (i = 0; i < sizeData; i++) + for (j = 0; j < sizeData; j++) + file >> b_h[i * sizeData + j]; + + //Allocate and copy QAP data from CPU memory to GPU global memory + GPUObject.memCopy(a_d, a_h, sizeData * sizeData); + GPUObject.memCopy(b_d, b_h, sizeData * sizeData); + + } + +public: + + ElemType* a_h; + ElemType* b_h; + ElemType* a_d; + ElemType* b_d; + +}; +#endif diff --git a/trunk/paradiseo-gpu/src/problems/eval/EvalOneMax.h b/trunk/paradiseo-gpu/src/problems/eval/EvalOneMax.h new file mode 100644 index 000000000..a177b447a --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/eval/EvalOneMax.h @@ -0,0 +1,79 @@ +/* + + 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 __EvalOneMax_H +#define __EvalOneMax_H + +/** + * Full Evaluation of the solution + */ + +template +class EvalOneMax: public eoEvalFunc { + + public: + + /** + * Constructor + */ + + EvalOneMax() { + } + + /** + * Destructor + */ + + ~EvalOneMax(void) { + } + + /** + * Full evaluation of the solution + * @param _bitVector the solution to evaluate + */ + + void operator()(EOT & _bitVector) { + + unsigned sum = 0; + + for (unsigned i = 0; i < _bitVector.size(); i++) + sum += _bitVector[i]; + + //set the solution fitness + _bitVector.fitness(sum); + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/problems/eval/OneMaxIncrEval.h b/trunk/paradiseo-gpu/src/problems/eval/OneMaxIncrEval.h new file mode 100644 index 000000000..7af4dc890 --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/eval/OneMaxIncrEval.h @@ -0,0 +1,91 @@ +/* + + 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 __OneMaxIncrEval_H +#define __OneMaxIncrEval_H + +#include + +/** + * Incremental Evaluation of OneMax + */ + +template +class OneMaxIncrEval: public moGPUEvalFunc { + +public: + + typedef typename Neighbor::EOT EOT; + typedef typename EOT::Fitness Fitness; + typedef typename EOT::ElemType T; + + /** + * Constructor + */ + + OneMaxIncrEval() { + } + + /** + * Destructor + */ + + ~OneMaxIncrEval() { + } + + /** + * Incremental evaluation of the OneMax solution(bit vector),function inline can be called from host or device + * @param _bitVector the solution to evaluate + * @param _fitness the fitness of the current solution + * @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 + */ + +inline __host__ __device__ Fitness operator() (T * _bitVector,Fitness _fitness, unsigned int * _index) { + + Fitness tmp=_fitness; + for(unsigned i=0;i + 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 __PPPEval_H +#define __PPPEval_H + +#include + +template +class PPPEval: public eoEvalFunc { + +public: + + /** + * Constructor + * @param _pppData the specific data problem useful to evalute solution( vector of 1 & _1 for PPP) + */ + + PPPEval(PPPData & _pppData) { + pppData = _pppData; + } + + /** + * Destructor + */ + + ~PPPEval() { + } + + /** + * Full evaluation of the solution + * @param _sol the solution to evaluate + */ + + void operator()(EOT & _sol) { + + int *H; + int tmp; + int tmp_1 = 0; + int tmp_2 = 0; + + H = new int[Nd]; + + for (int i = 0; i < Md; i++) { + tmp = 0; + for (int j = 0; j < Nd; j++) { + tmp += pppData.a_h[i * Nd + j] * _sol[j]; + } + + tmp_1 += abs(tmp) - tmp; + if (tmp > 0) + H[tmp-1]++; + } + + for (int j = 0; j < Nd; j++) { + tmp_2 += abs(pppData.H_h[j] - H[j]); + } + + _sol.fitness(ca * tmp_1 + cb * tmp_2); + + delete[] H; + + } + +protected: + + PPPData pppData; + +}; + +#endif + diff --git a/trunk/paradiseo-gpu/src/problems/eval/PPPIncrEval.h b/trunk/paradiseo-gpu/src/problems/eval/PPPIncrEval.h new file mode 100644 index 000000000..9ce5d22fc --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/eval/PPPIncrEval.h @@ -0,0 +1,109 @@ +/* + + 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 __PPPIncrEval_H +#define __PPPIncrEval_H + +#include + +/** + * Incremental Evaluation of PPP + */ + +template +class PPPIncrEval: public moGPUEvalFunc { + +public: + + typedef typename Neighbor::EOT EOT; + typedef typename EOT::Fitness Fitness; + typedef typename EOT::ElemType T; + + /** + * Constructor + */ + + PPPIncrEval() { + } + + /** + * Destructor + */ + + ~PPPIncrEval() { + } + + /** + * Incremental evaluation of the PPP solution,function inline can be called from host or device + * @param _sol the solution to evaluate + * @param _fitness the fitness of the current solution + * @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 + */ + + inline __host__ __device__ Fitness operator() (T* _sol,Fitness _fitness, unsigned int *_index) { + + int H[Nd]; + int S[Md]; + int tmp_1=0; + int tmp_2=0; + + for (unsigned i=0; i0) + H[S[i]-1]=H[S[i]-1]+1; + } + + for (unsigned j=0; j + 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 __QAPEval_H +#define __QAPEval_H + +#include + +template +class QAPEval: public eoEvalFunc { + +public: + + /** + * Constructor + * @param _qapData the specific data problem useful to evalute solution(flow & distance matrices of QAP problem) + */ + + QAPEval(QAPData & _qapData) { + qapData = _qapData; + } + + /** + * Destructor + */ + + ~QAPEval() { + } + + /** + * Full evaluation of the solution + * @param _sol the solution to evaluate + */ + + void operator()(EOT & _sol) { + int cost = 0; + unsigned int size = qapData.getSize(); + for (unsigned int i = 0; i < size; i++) + for (unsigned int j = 0; j < size; j++) { + cost += qapData.a_h[i * size + j] * qapData.b_h[_sol[i] * size + + _sol[j]]; + } + + _sol.fitness(cost); + } + +protected: + + QAPData qapData; + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/problems/eval/QAPIncrEval.h b/trunk/paradiseo-gpu/src/problems/eval/QAPIncrEval.h new file mode 100644 index 000000000..71a316e32 --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/eval/QAPIncrEval.h @@ -0,0 +1,122 @@ +/* + + 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 __QAPIncrEval_H +#define __QAPIncrEval_H + +#include + +/** + * Parallel Incremental Evaluation of QAP + */ + +template +class QAPIncrEval: public moGPUEvalFunc { + +public: + + typedef typename Neighbor::EOT EOT; + typedef typename EOT::Fitness Fitness; + typedef typename EOT::ElemType T; + + /** + * Constructor + */ + + QAPIncrEval() { + } + + /** + * Destructor + */ + + ~QAPIncrEval() { + } + + /** + * Incremental evaluation of the QAP solution,function inline can be called from host or device + * @param _sol the solution to evaluate + * @param _fitness the fitness of the current solution + * @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 + */ + +inline __host__ __device__ Fitness operator() (T * _sol,Fitness _fitness, unsigned int *_index) { + + Fitness tmp=_fitness; + + T tmp_sol[1]; + /* + * dev_a & dev_b are global device variable, data specific to QAP problem (flow & distance matices) + * _index[i] the first position of swap + * _index[i+1] the second position of swap + */ + for(unsigned i=0;i + 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 ue, + 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". + + 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 __PPPNeighbor_h +#define __PPPNeighbor_h + +#include +#include + +/** + * A GPU X-BitFlipping Neighbor + */ + +template +class PPPNeighbor: public moBackableNeighbor , public moXChangeNeighbor< + EOT> { +public: + + using moXChangeNeighbor::indices; + using moXChangeNeighbor::xChange; + using moXChangeNeighbor::key; + + /** + *Default Constructor + */ + + PPPNeighbor() : + moXChangeNeighbor () { + } + + /** + * Constructor + * @param _xFlip the number of bit to flip + */ + + PPPNeighbor(unsigned int _xFlip) : + moXChangeNeighbor (_xFlip) { + } + + /** + * Apply the K-Flip in solution + * @param _solution the solution to move + */ + + virtual void move(EOT& _solution) { + for (unsigned int i = 0; i < xChange; i++) + _solution[indices[i]] = -_solution[indices[i]]; + + _solution.invalidate(); + + } + + /** + * apply the K-Flip to restore the solution (use by moFullEvalByModif) + * @param _solution the solution to move back + */ + + virtual void moveBack(EOT& _solution) { + move(_solution); + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + + virtual std::string className() const { + return "PPPNeighbor"; + } + +}; + +#endif + diff --git a/trunk/paradiseo-gpu/src/problems/neighborhood/moGPUCustomizedNeighbor.h b/trunk/paradiseo-gpu/src/problems/neighborhood/moGPUCustomizedNeighbor.h new file mode 100644 index 000000000..65657a315 --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/neighborhood/moGPUCustomizedNeighbor.h @@ -0,0 +1,106 @@ +/* + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Jerémie Humeau, 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 _moGPUCustomizedNeighbor_h +#define _moGPUCustomizedNeighbor_h + +#include +#include +#include + +/** + * Neighbor related to a solution vector composed by two vectors + */ + +template +class moGPUCustomizedNeighbor: public moBackableNeighbor< moGPUSolType2Vector > , +public moXChangeNeighbor< moGPUSolType2Vector > { + +public: + + using moXChangeNeighbor< moGPUSolType2Vector >::indices; + using moXChangeNeighbor< moGPUSolType2Vector >::xChange; + using moXChangeNeighbor< moGPUSolType2Vector >::key; + /** + *Default Constructor + */ + + moGPUCustomizedNeighbor() : + moXChangeNeighbor< moGPUSolType2Vector > () { + } + + /** + * Constructor + * @param _xSwap the number of bit to swap + */ + + moGPUCustomizedNeighbor(unsigned int _xSwap) : + moXChangeNeighbor< moGPUSolType2Vector > (_xSwap) { + } + + /** + * move the solution + * @param _solution the solution to move + */ + + virtual void move(moGPUSolType2Vector & _solution) { + std::cout<<"_solution"< & _solution) { + move(_solution); + } + + /** + * Return the class name. + * @return the class name as a std::string + */ + virtual std::string className() const { + return "moGPUCustomizedNeighbor"; + } + +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/problems/types/PPPSolution.h b/trunk/paradiseo-gpu/src/problems/types/PPPSolution.h new file mode 100644 index 000000000..f04995c51 --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/types/PPPSolution.h @@ -0,0 +1,131 @@ +/* + + 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 __PPPSolution_H_ +#define __PPPSolution_H_ + +#include + +/** + * Implementation of PPP vector representation on GPU. + */ + +template + +class PPPSolution: public moGPUVector { + +public: + + using moGPUVector::vect; + using moGPUVector::N; + + /** + * Default constructor. + */ + + PPPSolution() : + moGPUVector () { + + } + + /** + *Constructor. + *@param _size The neighborhood size. + */ + + PPPSolution(unsigned _size) { + + N = _size; + + vect = new int[_size]; + + create(); + } + + /** + *Assignment operator + *@param _vector The vector passed to the function determine the new content. + *@return a new vector. + */ + + PPPSolution& operator=(const PPPSolution & _vector) { + + N = _vector.N; + vect = new int[N]; + for (unsigned i = 0; i < N; i++) + vect[i] = _vector.vect[i]; + fitness(_vector.fitness()); + return (*this); + + } + + /** + *Initializer of random PPP vector. + */ + void create() { + + for (int i = 0; i < N; i++) { + if ((rng.rand() % 2) == 0) + vect[i] = -1; + else + vect[i] = 1; + } + } + + /** + *Function inline to set the size of vector, called from host and device. + *@param _size the vector size + */ + + virtual void setSize(unsigned _size){ + N=_size; + } + /** + * Print the solution + */ + + virtual void printOn(std::ostream& os) const { + EO::printOn(os); + os << ' '; + os << N << ' '; + unsigned int i; + for (i = 0; i < N; i++) + os << vect[i] << ' '; + + } + +}; + +#endif + diff --git a/trunk/paradiseo-gpu/src/problems/types/moGPUCustomizeType.h b/trunk/paradiseo-gpu/src/problems/types/moGPUCustomizeType.h new file mode 100644 index 000000000..2e068badf --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/types/moGPUCustomizeType.h @@ -0,0 +1,116 @@ +/* + + 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 _moGPUCustomizeType_H_ +#define _moGPUCustomizeType_H_ + +/** + * Implementation of an Example of customized type + */ + +template +struct sol2Type { + + T1 tab1[SIZE]; + T2 tab2[SIZE]; + +inline __host__ __device__ sol2Type& operator=(const sol2Type _vector) { + for (unsigned i = 0; i < SIZE; i++) { + + tab1[i] = _vector.tab1[i]; + tab2[i] = _vector.tab2[i]; + } + return (*this); +} + +inline __host__ __device__ unsigned size() { + + return SIZE; + +} +}; + + +template +struct sol3Type { + + T1 tab1[SIZE]; + T2 tab2[SIZE]; + T3 tab3[SIZE]; + +inline __host__ __device__ sol3Type& operator=(const sol3Type _vector) { + + for (unsigned i = 0; i < SIZE; i++) { + + tab1[i] = _vector.tab1[i]; + tab2[i] = _vector.tab2[i]; + tab3[i] = _vector.tab3[i]; + } + return (*this); +} + +inline __host__ __device__ unsigned size() { + + return SIZE; + +} +}; + +template +struct sol4Type { + + T1 tab1[SIZE]; + T2 tab2[SIZE]; + T3 tab3[SIZE]; + T4 tab4[SIZE]; + +inline __host__ __device__ sol4Type& operator=(const sol4Type _vector) { + + for (unsigned i = 0; i < SIZE; i++) { + tab1[i] = _vector.tab1[i]; + tab2[i] = _vector.tab2[i]; + tab3[i] = _vector.tab3[i]; + tab4[i] = _vector.tab4[i]; + } + return (*this); +} + +inline __host__ __device__ unsigned size() { + + return SIZE; + +} +}; + +#endif diff --git a/trunk/paradiseo-gpu/src/problems/types/moGPUSolType2Vector.h b/trunk/paradiseo-gpu/src/problems/types/moGPUSolType2Vector.h new file mode 100755 index 000000000..3ed1eea58 --- /dev/null +++ b/trunk/paradiseo-gpu/src/problems/types/moGPUSolType2Vector.h @@ -0,0 +1,141 @@ +/* + + 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 _moGPUSolType2Vector_H_ +#define _moGPUSolType2Vector_H_ + +#include +#include + +/** + * An Example of a customized vector representation on GPU. + */ + +typedef struct sol2Type ElemType; +template +class moGPUSolType2Vector: public moGPUVector { + +public: + /** + * Define vector type of vector corresponding to Solution + */ + + using moGPUVector::vect; + using moGPUVector::N; + + /** + * Default constructor. + */ + + moGPUSolType2Vector() : + moGPUVector () { + } + + /** + *Constructor. + *@param _size The size of the vector to create. + */ + + moGPUSolType2Vector(unsigned _size) : + moGPUVector (_size) { + create(); + } + + /** + *Assignment operator + *@param _vector The vector passed to the function determine the new content. + *@return a new vector. + */ + + moGPUSolType2Vector & operator=( + const moGPUSolType2Vector & _vector) { + + vect[0] = _vector[0]; + if (!(_vector.invalid())) + fitness(_vector.fitness()); + else + (*this).invalidate(); + return (*this); + } + + /** + *How to fill the vector. + */ + + virtual void create() { + + for (int i = 0; i < vect[0].size(); i++) { + vect[0].tab1[i] = (int) (rng.rand() % (vect[0].size() - i) + i); + vect[0].tab2[i] = (float) (rng.rand() % (vect[0].size() - i) + i); + } + } + + /** + *Function inline to set the size of vector, called from host and device. + *@param _size the vector size + */ + + virtual void setSize(unsigned _size) { + N = _size; + } + /** + * Print the solution + */ + + virtual void printOn(std::ostream& os) const { + + EO::printOn(os); + os << ' '; + os << vect[0].size() << ' '; + unsigned int i; + for (i = 0; i < vect[0].size(); i++) { + os << vect[0].tab1[i] << ' '; + } + os << endl; + for (i = 0; i < vect[0].size(); i++) { + os << vect[0].tab2[i] << ' '; + } + os << endl; + + } + +inline __host__ __device__ unsigned size() { + + return N; + +} + +}; + +#endif