From 87dd3eaa92b81f8e319d5fa72065b3b5025af40c Mon Sep 17 00:00:00 2001 From: boufaras Date: Thu, 10 Feb 2011 13:43:14 +0000 Subject: [PATCH] git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2135 331e1502-861f-0410-8da2-ba01fb791d7f --- ParadisEO-GPU/src/eval/moCudaKswapEval.h | 337 +++++++++++------------ 1 file changed, 168 insertions(+), 169 deletions(-) diff --git a/ParadisEO-GPU/src/eval/moCudaKswapEval.h b/ParadisEO-GPU/src/eval/moCudaKswapEval.h index 39326faa6..80aab5148 100644 --- a/ParadisEO-GPU/src/eval/moCudaKswapEval.h +++ b/ParadisEO-GPU/src/eval/moCudaKswapEval.h @@ -1,36 +1,36 @@ /* - - Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 + + Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010 - Karima Boufaras, Thé Van LUONG + 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". + 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. + 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. + 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 - */ + ParadisEO WebSite : http://paradiseo.gforge.inria.fr + Contact: paradiseo-help@lists.gforge.inria.fr +*/ #ifndef moCudaKswapEval_H #define moCudaKswapEval_H @@ -42,172 +42,171 @@ */ template -class moCudaKswapEval: public moCudaEval { + class moCudaKswapEval: public moCudaEval { -public: + 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; + /** + * Define type of a solution corresponding to Neighbor + */ + typedef typename Neighbor::EOT EOT; + /** + * Define type of a vector corresponding to Solution + */ + typedef typename EOT::ElemType T; + /** + * Define type of a fitness corresponding to Solution + */ + typedef typename EOT::Fitness Fitness; - using moCudaEval::neighborhoodSize; - using moCudaEval::host_FitnessArray; - using moCudaEval::device_FitnessArray; - using moCudaEval::device_solution; - using moCudaEval::kernel_Dim; + using moCudaEval::neighborhoodSize; + using moCudaEval::host_FitnessArray; + using moCudaEval::device_FitnessArray; + using moCudaEval::device_solution; + using moCudaEval::kernel_Dim; - /** - * Constructor - * @param _neighborhoodSize the size of the neighborhood - * @param _incrEval the incremental evaluation - */ + /** + * Constructor + * @param _neighborhoodSize the size of the neighborhood + * @param _incrEval the incremental evaluation + */ - moCudaKswapEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) : - moCudaEval (_neighborhoodSize), incrEval(_incrEval) { - mutex = false; - mutex_kswap = false; + moCudaKswapEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) : + moCudaEval (_neighborhoodSize), incrEval(_incrEval) { + mutex = false; + mutex_kswap = false; + } + + /** + * Destructor + */ + ~moCudaKswapEval() { + if (mutex_kswap) { + cudaFree(&device_setSolution); + cudaFree(&device_tmp); + delete[] vect; + } + } + + /** + * Compute fitness for all solution neighbors in device without specific mapping + * @param _sol the solution which generate the neighborhood + */ + + virtual void neighborhoodEval(EOT & _sol) { + } + + /** + * Compute fitness for all solution neighbors in device with K-swap mapping + * @param _sol the solution which generate the neighborhood + * @param _mapping the array of mapping indexes for K-swap neighborhood + * @param _Kswap the number of swap + */ + + void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping, unsigned _Kswap) { + + // the solution size + unsigned _size = _sol.size(); + + // Get Current solution fitness + Fitness fitness = _sol.fitness(); + + //Case of Permutation + if (_Kswap == 1) { + if (!mutex) { + //Allocate the space for solution in the device global memory + cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); + mutex = true; } - /** - * Destructor - */ - ~moCudaKswapEval() { - if (mutex_kswap) { - cudaFree(&device_setSolution); - cudaFree(&device_tmp); - delete[] vect; - } + //Copy the solution vector from the host to device + cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T), + cudaMemcpyHostToDevice); + + //Launch the Kernel to compute all permutation neighbors fitness + kernelPermutation<<>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_size); + //Copy the result from device to host + cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize + * sizeof(Fitness), cudaMemcpyDeviceToHost); + } + //Case Kswap + else if (_Kswap > 1) { + if (!mutex_kswap) { + vect = new T[neighborhoodSize * _size]; + //Allocate the space for set of solution in the device global memory + cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize + * _size * sizeof(T)); + //Allocate the space to save temporary EOT element to swap + cudaMalloc((void**) &device_tmp.vect, neighborhoodSize + * sizeof(T)); + mutex_kswap = true; } - /** - * Compute fitness for all solution neighbors in device without specific mapping - * @param _sol the solution which generate the neighborhood - */ - - virtual void neighborhoodEval(EOT & _sol) { + for (int i = 0; i < neighborhoodSize; i++) { + for (int j = 0; j < _size; j++) { + vect[j + i * _size] = _sol.vect[j]; + } } - /** - * Compute fitness for all solution neighbors in device with K-swap mapping - * @param _sol the solution which generate the neighborhood - * @param _mapping the array of mapping indexes for K-swap neighborhood - * @param _Kswap the number of swap - */ + //Copy the set of solution from the host to device + cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size + * sizeof(T), cudaMemcpyHostToDevice); - void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping, unsigned _Kswap) { + //Launch the Kernel to compute all Kswap neighbors fitness + kernelKswap<<>>(incrEval,device_setSolution,device_tmp,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap,_size); - // the solution size - unsigned _size = _sol.size(); + //Copy the result from device to host + cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize + * sizeof(Fitness), cudaMemcpyDeviceToHost); - // Get Current solution fitness - Fitness fitness = _sol.fitness(); + } + } - //Case of Permutation - if (_Kswap == 1) { - if (!mutex) { - //Allocate the space for solution in the device global memory - cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); - mutex = true; - } + /** + * Compute fitness for all solution neighbors(K-flip of binary solution) in device + * @param _sol the solution which generate the neighborhood + * @param _mapping the array of mapping indexes for k-flip neighborhood + * @param _Kflip the number of flip to do + */ - //Copy the solution vector from the host to device - cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T), - cudaMemcpyHostToDevice); + void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping, + unsigned _Kflip) { - //Launch the Kernel to compute all permutation neighbors fitness - kernelPermutation<<>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_size); - //Copy the result from device to host - cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize - * sizeof(Fitness), cudaMemcpyDeviceToHost); - } - //Case Kswap - else if (_Kswap > 1) { - if (!mutex_kswap) { - vect = new T[neighborhoodSize * _size]; - //Allocate the space for set of solution in the device global memory - cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize - * _size * sizeof(T)); - //Allocate the space to save temporary EOT element to swap - cudaMalloc((void**) &device_tmp.vect, neighborhoodSize - * sizeof(T)); - mutex_kswap = true; - } + // the solution size + unsigned _size = _sol.size(); - for (int i = 0; i < neighborhoodSize; i++) { - for (int j = 0; j < _size; j++) { - vect[j + i * _size] = _sol.vect[j]; - } - } + // Get Current solution fitness + Fitness fitness = _sol.fitness(); + if (!mutex) { + //Allocate the space for solution in the device global memory + cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); + mutex = true; + } - //Copy the set of solution from the host to device - cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size - * sizeof(T), cudaMemcpyHostToDevice); + //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 Kswap neighbors fitness - kernelKswap<<>>(incrEval,device_setSolution,device_tmp,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap,_size); + //Launch the Kernel to compute all flip neighbors fitness + kernelKflip<<>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kflip); - //Copy the result from device to host - cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize - * sizeof(Fitness), cudaMemcpyDeviceToHost); + //Copy the result from device to host + cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize + * sizeof(Fitness), cudaMemcpyDeviceToHost); + } - } - } + protected: - /** - * Compute fitness for all solution neighbors(K-flip of binary solution) in device - * @param _sol the solution which generate the neighborhood - * @param _mapping the array of mapping indexes for k-flip neighborhood - * @param _Kflip the number of flip to do - */ - - void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping, - unsigned _Kflip) { - - // the solution size - unsigned _size = _sol.size(); - - // Get Current solution fitness - Fitness fitness = _sol.fitness(); - if (!mutex) { - //Allocate the space for solution in the device global memory - cudaMalloc((void**) &device_solution.vect, _size * sizeof(T)); - mutex = true; - } - - //Copy the solution vector from the host to device - cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T), - cudaMemcpyHostToDevice); - - //Launch the Kernel to compute all flip neighbors fitness - kernelKflip<<>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kflip); - - //Copy the result from device to host - cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize - * sizeof(Fitness), cudaMemcpyDeviceToHost); - } - - protected: - - IncrementEval & incrEval; - //NeighborhoodSize copy of solution - EOT device_setSolution; - //NeighborhoodSize element of EOT - EOT device_tmp; - //Vector of neighborhoodSize copy of solution - T * vect; - bool mutex_kswap; - bool mutex; + IncrementEval & incrEval; + //NeighborhoodSize copy of solution + EOT device_setSolution; + //NeighborhoodSize element of EOT + EOT device_tmp; + //Vector of neighborhoodSize copy of solution + T * vect; + bool mutex_kswap; + bool mutex; }; #endif -