GPU -> GPU-good

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

ParadisEO-GPU-good/src/cudaType/moCudaBitVector.h

@@ -0,0 +1,107 @@
+/*
+  <moCudaBitVector.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Karima Boufaras, Thé Van LUONG
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaBitVector_H_
+#define _moCudaBitVector_H_
+
+#include <cudaType/moCudaVector.h>
+
+/**
+ * Implementation of Bit vector representation on CUDA.
+ */
+
+template<class Fitness>
+
+class moCudaBitVector: public moCudaVector<bool, Fitness> {
+
+ public:
+
+  using moCudaVector<bool, Fitness>::vect;
+  using moCudaVector<bool, Fitness>::N;
+
+  /**
+   * Default constructor.
+   */
+
+ moCudaBitVector() :
+  moCudaVector<bool, Fitness> () {
+
+  }
+
+  /**
+   *Constructor.
+   *@param _size The neighborhood size.
+   */
+
+  moCudaBitVector(unsigned _size) {
+
+    N = _size;
+
+    vect = new bool[_size];
+
+    create();
+  }
+
+  /**
+   *Constructor.
+   *@param _size The neighborhood size.
+   *@param _b Value to assign to vector.
+   */
+
+  moCudaBitVector(unsigned _size, bool _b) {
+
+    N = _size;
+
+    vect = new bool[_size];
+
+    for (unsigned i = 0; i < _size; i++)
+      vect[i] = _b;
+  }
+
+  /**
+   *Initializer of random bit vector.
+   */
+
+  void create() {
+
+    for (unsigned i = 0; i < N; i++) {
+
+      vect[i] = (bool) round((float) rng.rand() / RAND_MAX);
+
+    }
+  }
+
+};
+
+#endif

ParadisEO-GPU-good/src/cudaType/moCudaIntVector.h

@@ -0,0 +1,113 @@
+/*
+  <moCudaIntVector.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Karima Boufaras, Thé Van LUONG
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaIntVector_H_
+#define _moCudaIntVector_H_
+
+#include <cudaType/moCudaVector.h>
+
+/**
+ * Implementation of integer vector representation on CUDA.
+ */
+
+template<class Fitness>
+
+class moCudaIntVector: public moCudaVector<int, Fitness> {
+
+ public:
+
+  using moCudaVector<int, Fitness>::vect;
+  using moCudaVector<int, Fitness>::N;
+
+  /**
+   * Default constructor.
+   */
+
+ moCudaIntVector() :
+  moCudaVector<int, Fitness> () {
+
+  }
+
+  /**
+   *Constructor.
+   *@param _size The neighborhood size.
+   */
+
+  moCudaIntVector(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.
+   */
+
+  moCudaIntVector& operator=(const moCudaIntVector & _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 integer vector.
+   */
+  void create() {
+
+    unsigned random;
+    int temp;
+    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;
+      temp = vect[i];
+      vect[i] = vect[random];
+      vect[random] = temp;
+    }
+  }
+
+};
+
+#endif

ParadisEO-GPU-good/src/cudaType/moCudaVector.h

@@ -0,0 +1,159 @@
+/*
+  <moCudaVector.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Thé Van LUONG, Karima Boufaras
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaVector_H_
+#define _moCudaVector_H_
+
+#include <eo>
+
+/**
+ * Implementation of vector representation on CUDA.
+ */
+
+template<class ElemT, class Fitness>
+
+  class moCudaVector: public EO<Fitness> {
+
+ public:
+
+    typedef ElemT ElemType;
+    ElemType * vect;
+    unsigned N;
+
+    /**
+     * Default constructor.
+     */
+
+ moCudaVector() :
+    N(0) {
+    }
+
+    /**
+     *Constructor.
+     *@param _size The neighborhood size.
+     */
+
+    moCudaVector(unsigned _size) {
+
+      N = _size;
+
+      vect = new ElemType[_size];
+
+      create();
+    }
+
+    /**
+     * Destructor.
+     */
+
+    ~moCudaVector() {
+      if (N > 1)
+	delete[] vect;
+    }
+
+    /**
+     *How to fill the vector.
+     */
+
+    virtual void create() {
+
+    }
+
+    /**
+     *Assignment operator
+     *@param _vector The vector passed to the function determine the new content.
+     *@return a new vector.
+     */
+
+    virtual moCudaVector& operator=(const moCudaVector & _vector) {
+
+      N = _vector.N;
+      vect = new ElemType[N];
+      for (unsigned i = 0; i < N; i++)
+	vect[i] = _vector.vect[i];
+      fitness(_vector.fitness());
+      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;
+
+    }
+
+    virtual void printOn(std::ostream& os) const {
+      EO<Fitness>::printOn(os);
+      os << ' ';
+      os << N << ' ';
+      unsigned int i;
+      for (i = 0; i < N; i++)
+	os << vect[i] << ' ';
+
+    }
+
+};
+
+#endif

ParadisEO-GPU-good/src/eval/moCudaEval.h

@@ -0,0 +1,136 @@
+/*
+  <moCudaEval.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Karima Boufaras, Thé Van LUONG
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef moCudaEval_H
+#define moCudaEval_H
+#include <eval/moEval.h>
+
+/**
+ * Abstract class for evaluation on GPU
+ */
+template<class Neighbor>
+class moCudaEval: public moEval<Neighbor> {
+
+ public:
+
+  /**
+   * Define type of a solution corresponding to Neighbor
+   */
+  typedef typename Neighbor::EOT EOT;
+
+  typedef typename EOT::Fitness Fitness;
+
+  /**
+   * Constructor
+   * @param _neighborhoodSize the size of the neighborhood
+   */
+
+  moCudaEval(unsigned int _neighborhoodSize) {
+
+    neighborhoodSize = _neighborhoodSize;
+    host_FitnessArray = new Fitness[neighborhoodSize];
+    cudaMalloc((void**) &device_FitnessArray, neighborhoodSize
+	       * sizeof(Fitness));
+    kernel_Dim = neighborhoodSize / BLOCK_SIZE + ((neighborhoodSize
+						   % BLOCK_SIZE == 0) ? 0 : 1);
+  }
+
+  /**
+   * Destructor
+   */
+
+  ~moCudaEval() {
+
+    delete[] host_FitnessArray;
+    cudaFree(device_FitnessArray);
+    cudaFree(&device_solution);
+  }
+
+  /**
+   * Set fitness of a solution neighbors
+   *@param _sol the solution which generate the neighborhood
+   *@param _neighbor the current neighbor
+   */
+
+  void operator()(EOT & _sol, Neighbor & _neighbor) {
+
+    _neighbor.fitness(host_FitnessArray[_neighbor.index()]);
+
+  }
+
+  /**
+   * Compute fitness for all  Kswap solution neighbors in device
+   * @param _sol the solution which generate the neighborhood
+   * @param _mapping the neighborhood mapping
+   * @param _Kswap the number of swap
+   */
+  virtual void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping,
+				     unsigned _Kswap) {
+
+  }
+
+  /**
+   * Compute fitness for all Kflip solution neighbors in device
+   * @param _sol the solution which generate the neighborhood
+   * @param _mapping the neighborhood mapping
+   * @param _Kflip the number of bit to flip
+   */
+  virtual void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping,
+				     unsigned _Kflip) {
+
+  }
+
+  /**
+   * Compute fitness for all solution neighbors in device
+   * @param _sol the solution which generate the neighborhood
+   */
+
+  virtual void neighborhoodEval(EOT & _sol)=0;
+
+ protected:
+
+  //the host array to save all neighbors fitness
+  Fitness * host_FitnessArray;
+  //the device array to save neighbors fitness computed in device
+  Fitness * device_FitnessArray;
+  //the device solution
+  EOT device_solution;
+  //the size of neighborhood
+  unsigned int neighborhoodSize;
+  //Cuda kernel dimension
+  int kernel_Dim;
+
+};
+
+#endif

ParadisEO-GPU-good/src/eval/moCudaEvalFunc.h

@@ -0,0 +1,84 @@
+/*
+  <moCudaEvalFunc.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Thé Van LUONG, Karima Boufaras
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef __moCudaEvalFunc_H
+#define __moCudaEvalFunc_H
+
+/**
+ * Abstract class for CUDA evaluation of neighbor
+ */
+
+template<class Neighbor>
+class moCudaEvalFunc {
+
+ public:
+
+  /**
+   * Define type of a solution corresponding to Neighbor
+   */
+
+  typedef typename Neighbor::EOT EOT;
+  typedef typename EOT::ElemType T;
+  typedef typename EOT::Fitness Fitness;
+
+  /**
+   * Constructor
+   */
+
+  moCudaEvalFunc() {
+  }
+
+  /**
+   * Destructor
+   */
+
+  ~moCudaEvalFunc() {
+  }
+
+  /**
+   *Virtual functor to compute fitness of a solution neighbor
+   *@param _solution the solution which generate the neighborhood
+   *@param _fitness the current solution fitness
+   *@param _index the index neighbor
+   */
+
+  virtual inline __host__ __device__ Fitness operator() (EOT & _solution,Fitness _fitness, unsigned int * _index){
+
+    return _fitness;
+
+  };
+
+
+};
+#endif

ParadisEO-GPU-good/src/eval/moCudaKswapEval.h

@@ -0,0 +1,212 @@
+/*
+  <moCudaKswapEval.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Karima Boufaras, Thé Van LUONG
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef moCudaKswapEval_H
+#define moCudaKswapEval_H
+#include <eval/moCudaEval.h>
+#include <eval/moCudakernelEval.h>
+
+/**
+ * class for the K-swap neighborhood evaluation
+ */
+
+template<class Neighbor, class IncrementEval>
+  class moCudaKswapEval: public moCudaEval<Neighbor> {
+
+ public:
+
+    /**
+     * Define type of a solution corresponding to Neighbor
+     */
+    typedef typename Neighbor::EOT EOT;
+    /**
+     * Define type of a vector corresponding to Solution
+     */
+    typedef typename EOT::ElemType T;
+    /**
+     * Define type of a fitness corresponding to Solution
+     */
+    typedef typename EOT::Fitness Fitness;
+
+    using moCudaEval<Neighbor>::neighborhoodSize;
+    using moCudaEval<Neighbor>::host_FitnessArray;
+    using moCudaEval<Neighbor>::device_FitnessArray;
+    using moCudaEval<Neighbor>::device_solution;
+    using moCudaEval<Neighbor>::kernel_Dim;
+
+    /**
+     * Constructor
+     * @param _neighborhoodSize the size of the neighborhood
+     * @param _incrEval the incremental evaluation
+     */
+
+ moCudaKswapEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) :
+    moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
+      mutex = false;
+      mutex_kswap = false;
+    }
+
+    /**
+     * Destructor
+     */
+    ~moCudaKswapEval() {
+      if (mutex_kswap) {
+	cudaFree(&device_setSolution);
+	cudaFree(&device_tmp);
+	delete[] vect;
+      }
+    }
+
+    /**
+     * Compute fitness for all solution neighbors in device without specific mapping
+     * @param _sol the solution which generate the neighborhood
+     */
+
+    virtual void neighborhoodEval(EOT & _sol) {
+    }
+
+    /**
+     * Compute fitness for all solution neighbors in device with K-swap mapping
+     * @param _sol the solution which generate the neighborhood
+     * @param _mapping the array of mapping indexes for K-swap neighborhood
+     * @param _Kswap the number of swap
+     */
+
+    void neighborhoodKswapEval(EOT & _sol, unsigned * _mapping, unsigned _Kswap) {
+
+      // the solution size
+      unsigned _size = _sol.size();
+
+      // Get Current solution fitness
+      Fitness fitness = _sol.fitness();
+
+      //Case of Permutation
+      if (_Kswap == 1) {
+	if (!mutex) {
+	  //Allocate the space for solution in the device global memory
+	  cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
+	  mutex = true;
+	}
+
+	//Copy the solution vector from the host to device
+	cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
+		   cudaMemcpyHostToDevice);
+
+	//Launch the Kernel to compute all  permutation neighbors fitness
+	kernelPermutation<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_size);
+	//Copy the result from device to host
+	cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
+		   * sizeof(Fitness), cudaMemcpyDeviceToHost);
+      }
+      //Case Kswap
+      else if (_Kswap > 1) {
+	if (!mutex_kswap) {
+	  vect = new T[neighborhoodSize * _size];
+	  //Allocate the space for set of solution in the device global memory
+	  cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize
+		     * _size * sizeof(T));
+	  //Allocate the space to save temporary EOT element to swap
+	  cudaMalloc((void**) &device_tmp.vect, neighborhoodSize
+		     * sizeof(T));
+	  mutex_kswap = true;
+	}
+
+	for (int i = 0; i < neighborhoodSize; i++) {
+	  for (int j = 0; j < _size; j++) {
+	    vect[j + i * _size] = _sol.vect[j];
+	  }
+	}
+
+	//Copy the set of solution from the host to device
+	cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size
+		   * sizeof(T), cudaMemcpyHostToDevice);
+
+	//Launch the Kernel to compute all  Kswap neighbors fitness
+	kernelKswap<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_setSolution,device_tmp,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kswap,_size);
+
+	//Copy the result from device to host
+	cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
+		   * sizeof(Fitness), cudaMemcpyDeviceToHost);
+
+      }
+    }
+
+    /**
+     * Compute fitness for all solution neighbors(K-flip of binary solution) in device
+     * @param _sol the solution which generate the neighborhood
+     * @param _mapping the array of mapping indexes for k-flip neighborhood
+     * @param _Kflip the number of flip to do
+     */
+
+    void neighborhoodKflipEval(EOT & _sol, unsigned * _mapping,
+			       unsigned _Kflip) {
+
+      // the solution  size
+      unsigned _size = _sol.size();
+
+      // Get Current solution fitness
+      Fitness fitness = _sol.fitness();
+      if (!mutex) {
+	//Allocate the space for solution in the device global memory
+	cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
+	mutex = true;
+      }
+
+      //Copy the solution vector from the host to device
+      cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
+		 cudaMemcpyHostToDevice);
+
+      //Launch the Kernel to compute all flip neighbors fitness
+      kernelKflip<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize,_mapping,_Kflip);
+
+      //Copy the result from device to host
+      cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
+		 * sizeof(Fitness), cudaMemcpyDeviceToHost);
+    }
+
+ protected:
+
+    IncrementEval & incrEval;
+    //NeighborhoodSize copy of solution
+    EOT device_setSolution;
+    //NeighborhoodSize element of EOT
+    EOT device_tmp;
+    //Vector of neighborhoodSize copy of solution
+    T * vect;
+    bool mutex_kswap;
+    bool mutex;
+
+};
+
+#endif

ParadisEO-GPU-good/src/eval/moCudaVectorEval.h

@@ -0,0 +1,113 @@
+/*
+ <moCudaVectorEval.h>
+ Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+ Karima Boufaras, Thé Van LUONG
+
+ This software is governed by the CeCILL license under French law and
+ abiding by the rules of distribution of free software.  You can  use,
+ modify and/ or redistribute the software under the terms of the CeCILL
+ license as circulated by CEA, CNRS and INRIA at the following URL
+ "http://www.cecill.info".
+
+ As a counterpart to the access to the source code and  rights to copy,
+ modify and redistribute granted by the license, users are provided only
+ with a limited warranty  and the software's author,  the holder of the
+ economic rights,  and the successive licensors  have only  limited liability.
+
+ In this respect, the user's attention is drawn to the risks associated
+ with loading,  using,  modifying and/or developing or reproducing the
+ software by the user in light of its specific status of free software,
+ that may mean  that it is complicated to manipulate,  and  that  also
+ therefore means  that it is reserved for developers  and  experienced
+ professionals having in-depth computer knowledge. Users are therefore
+ encouraged to load and test the software's suitability as regards their
+ requirements in conditions enabling the security of their systems and/or
+ data to be ensured and,  more generally, to use and operate it in the
+ same conditions as regards security.
+ The fact that you are presently reading this means that you have had
+ knowledge of the CeCILL license and that you accept its terms.
+
+ ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+ Contact: paradiseo-help@lists.gforge.inria.fr
+ */
+
+#ifndef moCudaVectorEval_H
+#define moCudavectorEval_H
+#include <eval/moCudaEval.h>
+#include <eval/moCudakernelEval.h>
+
+/**
+ * class for the parallel evaluation of neighborhood without mapping
+ */
+
+template<class Neighbor, class IncrementEval>
+class moCudaVectorEval: public moCudaEval<Neighbor> {
+
+public:
+
+	/**
+	 * Define type of a solution corresponding to Neighbor
+	 */
+	typedef typename Neighbor::EOT EOT;
+	/**
+	 * Define type of a vector corresponding to Solution
+	 */
+	typedef typename EOT::ElemType T;
+	/**
+	 * Define type of a fitness corresponding to Solution
+	 */
+	typedef typename EOT::Fitness Fitness;
+
+	using moCudaEval<Neighbor>::neighborhoodSize;
+	using moCudaEval<Neighbor>::host_FitnessArray;
+	using moCudaEval<Neighbor>::device_FitnessArray;
+	using moCudaEval<Neighbor>::device_solution;
+	using moCudaEval<Neighbor>::kernel_Dim;
+
+	/**
+	 * Constructor
+	 * @param _neighborhoodSize the size of the neighborhood
+	 * @param _incrEval the incremental evaluation
+	 */
+
+	moCudaVectorEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) :
+			moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
+		}
+	/**
+	 * Compute fitness for all solution neighbors in device
+	 * @param _sol the solution which generate the neighborhood
+	 */
+
+	virtual void neighborhoodEval(EOT & _sol) {
+
+		// the solution vector size
+		unsigned _size = _sol.size();
+
+		// Get Current solution fitness
+		Fitness fitness = _sol.fitness();
+
+		//Allocate the space for solution in the  global memory of device
+		cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
+
+		//Copy the solution vector from the host to device
+		cudaMemcpy(device_solution.vect, _sol.vect, _size * sizeof(T),
+				cudaMemcpyHostToDevice);
+
+		//Launch the Kernel to compute all neighbors fitness
+		kernelEval<EOT,Fitness,Neighbor,IncrementEval><<<kernel_Dim,BLOCK_SIZE >>>(incrEval,device_solution,device_FitnessArray,fitness,neighborhoodSize);
+
+		//Copy the result from device to host
+		cudaMemcpy(host_FitnessArray, device_FitnessArray, neighborhoodSize
+				* sizeof(Fitness), cudaMemcpyDeviceToHost);
+
+	}
+
+protected:
+
+	IncrementEval & incrEval;
+
+};
+
+#endif
+

ParadisEO-GPU-good/src/eval/moCudakernelEval.h

@@ -0,0 +1,206 @@
+/*
+  <moCudakernelEval.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Karima Boufaras, Thé Van LUONG
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+#ifndef __moCudakernelEval_H
+#define __moCudakernelEval_H
+
+
+///////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
+/**
+ * The kernel function called from the host and executed in device to compute all neighbors fitness at one time
+ * with linear mapping
+ * @param _eval how to evaluate each neighbor
+ * @param _solution the representation of solution( vector of int,float....)
+ * @param _allFitness the array of Fitness to save all neighbors fitness
+ * @param _fitness the current solution fitness
+ * @param _neighborhoodsize the size of the neighborhood
+ */
+
+template<class EOT, class Fitness, class Neighbor, class IncrementEval>
+
+  __global__ void kernelEval(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
+			     Fitness _fitness, unsigned _neighborhoodsize) {
+
+  // The thread identifier within a grid block's
+  int id = blockIdx.x * blockDim.x + threadIdx.x;
+  // array to save index to be changed
+  unsigned int index[1];
+  // In this representation each id identify one and only one neighbor in neighborhood
+  if (id < _neighborhoodsize) {
+    //Change the id'th element of solution
+    index[0]=id;
+    //Compute fitness for id'th neighbor
+    _allFitness[id] = _eval(_solution, _fitness,index);
+  }
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
+/**
+ * The kernel function called from the host and executed in device to compute all flip neighbors fitness at one time
+ * @param _eval how to evaluate each neighbor
+ * @param _solution representation of solution to flip
+ * @param _allFitness Array of Fitness type to save all neighbors fitness
+ * @param _fitness the current solution fitness
+ * @param _neighborhoodsize the size of the neighborhood
+ * @param _mapping the neighborhood mapping
+ * @param _Kflip the number of bit to flip
+ */
+
+template<class EOT, class Fitness, class Neighbor, class IncrementEval>
+
+  __global__ void kernelKflip(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
+			      Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _Kflip) {
+
+  // The thread identifier within a grid block's
+  int id = blockIdx.x * blockDim.x + threadIdx.x;
+  //save temporary fitness
+  unsigned tmp_fitness;
+  //counter of number of flip to do
+  unsigned i;
+  // array to save index to be changed
+  unsigned index[1];
+  // In this representation each id identify one and only one neighbor in neighborhood
+  if (id < _neighborhoodsize) {
+    //Init fitness with fitness of solution
+    tmp_fitness=_fitness;
+    //Evaluate neighbor after Kflip
+    for(i=0;i<=_Kflip;i++){
+      //The designed index to flip
+      index[0]=_mapping[id + i * _neighborhoodsize];
+      //Evaluate the neighbor
+      tmp_fitness= _eval(_solution, tmp_fitness, index);
+
+    }
+    //The final  fitness of the Id'th neighbor
+    _allFitness[id]=tmp_fitness;
+
+  }
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
+/**
+ * The kernel function called from the host and executed in device to compute all swap neighbors fitness at one time
+ * @param _eval how to evaluate each neighbor
+ * @param _solution representation ofsolution to swap
+ * @param _sol_tmp to save temporary a solution element to swap
+ * @param _allFitness Array of Fitness type to save all neighbors fitness
+ * @param _fitness the current solution fitness
+ * @param _neighborhoodsize the size of the neighborhood
+ * @param _mapping the neighborhood mapping
+ * @param _Kswap the number of swap to do
+ * @param _size the solution size
+ */
+
+template<class EOT,class Fitness, class Neighbor, class IncrementEval>
+
+  __global__ void kernelKswap(IncrementEval _eval,EOT _solution ,EOT _sol_tmp, Fitness* _allFitness,
+			      Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _Kswap,unsigned _size) {
+
+  // The thread identifier within a grid block's
+  int id = blockIdx.x * blockDim.x + threadIdx.x;
+  //save temporary fitness
+  int tmp_fitness;
+  //counter of number of swap to do
+  unsigned i;
+  // array to save index to be changed, solution size & thread id
+  unsigned index[4];
+  // In this representation each id identify one and only one neighbor in neighborhood
+  if (id < _neighborhoodsize) {
+    //the first index to swap
+    index[0]=_mapping[id];
+    //the second index to swap
+    index[1]=_mapping[id +_neighborhoodsize];
+    //the solution size
+    index[2]=_size;
+    //the thread id
+    index[3]=id;
+    //Init the temporary fitness with the initial solution fitness
+    tmp_fitness=_fitness;
+
+    //Evaluate neighbor after K-swap
+    for(i=2;i<=_Kswap+1;i++){
+      //Evaluate neighbor with index case
+      tmp_fitness=_eval(_solution, tmp_fitness, index);
+      //Permut the solution
+      _sol_tmp[id]=_solution[index[0]+id*index[2]];
+      _solution[index[0]+id*index[2]]=_solution[index[1]+id*index[2]];
+      _solution[index[1]+id*index[2]]=_sol_tmp[id];
+      //Init the next swap to do
+      index[0]=index[1];
+      index[1]=_mapping[id +i*_neighborhoodsize];
+
+    }
+    //save the final fitness of the id'th neighbor
+    _allFitness[id]=tmp_fitness;
+  }
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
+/**
+ * The kernel function called from the host and executed in device to compute all permutation neighbors fitness at one time
+ * @param _eval how to evaluate each neighbor
+ * @param _solution representation of  solution
+ * @param _allFitness Array of Fitness type to save all neighbors fitness
+ * @param _fitness the current solution fitness
+ * @param _neighborhoodsize the size of the neighborhood
+ * @param _mapping the neighborhood mapping
+ * @param _size the solution size
+ */
+
+template<class EOT, class Fitness, class Neighbor, class IncrementEval>
+  __global__ void kernelPermutation(IncrementEval _eval, EOT _solution, Fitness* _allFitness,
+				    Fitness _fitness, unsigned _neighborhoodsize, unsigned * _mapping,unsigned _size) {
+
+  // The thread identifier within a grid block's
+  int id = blockIdx.x * blockDim.x + threadIdx.x;
+  // array to save index to be changed, solution size
+  unsigned index[4];
+  // In this representation each id identify one and only one neighbor in neighborhood
+  if (id < _neighborhoodsize) {
+    //The first index of permutation
+    index[0]=_mapping[id];
+    //The second index of permutation
+    index[1]=_mapping[id +_neighborhoodsize];
+    //The solution size
+    index[2]=_size;
+    //Puch 0 in the 3 index
+    index[3]=0;
+    _allFitness[id]=_eval(_solution,_fitness,index);
+  }
+}
+
+#endif

ParadisEO-GPU-good/src/memory/moCudaAllocator.h

@@ -0,0 +1,81 @@
+/*
+  <moCudaAllocator.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaAllocator_H_
+#define _moCudaAllocator_H_
+
+// CUDA includes
+#include <cutil.h>
+
+/**
+ *  class for allocation data on GPU global memory
+ */
+
+class moCudaAllocator {
+
+ public:
+
+  /**
+   * Constructor
+   */
+
+  moCudaAllocator() {
+  }
+
+  /**
+   *Allocate data on GPU global memory
+   *@param _data the data to allocate on GPU global memory
+   *@param _dataSize the size of data to allocate on GPU memory
+   */
+  template<typename T>
+    void operator()(T* & _data, unsigned _dataSize) {
+
+    //Allocate data in GPU memory
+    CUDA_SAFE_CALL(cudaMalloc((void**) &_data, _dataSize * sizeof(T)));
+
+    // Check if data allocation is failed
+    CUT_CHECK_ERROR("Allocation of data on GPU global memory failed");
+
+  }
+
+  /**
+   * Destructor
+   */
+
+  ~moCudaAllocator() {
+  }
+
+};
+
+#endif

ParadisEO-GPU-good/src/memory/moCudaCopy.h

@@ -0,0 +1,117 @@
+/*
+  <moCudaCopy.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaCopy_H_
+#define _moCudaCopy_H_
+
+// CUDA includes
+#include <cutil.h>
+
+/**
+ *  class to copy data  from CPU memory to  GPU global memory and vice versa
+ */
+
+class moCudaCopy {
+
+public:
+
+  /**
+   * Constructor
+   */
+
+  moCudaCopy() {
+  }
+
+  /**
+   *Copy data from CPU memory to GPU global memory (default copy)
+   *@param _data the data representation where the data will be copied
+   *@param _dataTocpy the data to copy from CPU memory to GPU memory
+   *@param _dataSize the size of data to copy 
+   */
+  template<typename T>
+    void operator()(T* & _data, T * & _dataTocpy, unsigned _dataSize) {
+
+    //copy data from CPU memory to GPU memory
+    CUDA_SAFE_CALL(cudaMemcpy(_data, _dataTocpy, _dataSize * sizeof(T),
+			      cudaMemcpyHostToDevice));
+
+    // Check if the copy of data is failed
+    CUT_CHECK_ERROR("Copy of data from CPU to GPU global memory failed");
+  }
+
+  /**
+   *Copy device data from GPU global memory to global variable declared in device
+   *@param _dev_data the device global variable
+   *@param _dataTocpy the data to copy GPU global memory to GPU global variable
+   */
+  template<typename T>
+    void operator()(T* & _dev_data, T * & _dataTocpy) {
+
+    //Copy n bytes from the memory area pointed to by _dataTocpy to the memory area pointed to by offset bytes from the start of symbol _dev_data
+
+    cudaMemcpyToSymbol(_dev_data, &_dataTocpy, sizeof(_dataTocpy));
+
+  }
+
+  /**
+   *Copy data from CPU memory to GPU global memory and vice versa
+   *@param _data the data representation where the data will be copied
+   *@param _dataTocpy the data to copy from CPU memory to GPU memory and vice versa
+   *@param _dataSize the size of data to copy 
+   *@param _HostToDevice the direction of copy(true if copy will be done from CPU memory to GPU memory)
+   */
+  template<typename T>
+    void operator()(T* & _data, T * & _dataTocpy, unsigned _dataSize,
+		    bool _HostToDevice) {
+
+    if (_HostToDevice) {
+
+      //copy data from CPU memory to GPU global memory
+      CUDA_SAFE_CALL(cudaMemcpy(_data, _dataTocpy, _dataSize * sizeof(T),
+				cudaMemcpyHostToDevice));
+    }
+
+    else {
+      //copy data from GPU global memory to GPU memory
+      CUDA_SAFE_CALL(cudaMemcpy(_data, _dataTocpy, _dataSize * sizeof(T),
+				cudaMemcpyDeviceToHost));
+    }
+
+  }
+
+  ~moCudaCopy() {
+  }
+
+};
+#endif

ParadisEO-GPU-good/src/memory/moCudaDesallocator.h

@@ -0,0 +1,77 @@
+/*
+  <moCudaDesallocator.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaDesallocator_H_
+#define _moCudaDesallocator_H_
+
+// CUDA includes
+#include <cutil.h>
+
+/**
+ *  class for Desallocation of data from GPU global memory
+ */
+
+class moCudaDesallocator {
+
+ public:
+
+  /**
+   * Constructor
+   */
+
+  moCudaDesallocator() {
+  }
+
+  /**
+   *Desallocate data on GPU global memory
+   *@param _data the data to desallocate from GPU global memory
+   */
+  template<typename T>
+    void operator()(T* & _data) {
+
+    //Desallocate data from GPU global memory
+    cudaFree(_data);
+
+  }
+
+  /**
+   * Destructor
+   */
+
+  ~moCudaDesallocator() {
+  }
+
+};
+
+#endif

ParadisEO-GPU-good/src/memory/moCudaObject.h

@@ -0,0 +1,91 @@
+/*
+  <moCudaObject.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaObject_H_
+#define _moCudaObject_H_
+
+// CUDA includes
+#include <cutil.h>
+#include <memory/moCudaAllocator.h>
+#include <memory/moCudaDesallocator.h>
+#include <memory/moCudaCopy.h>
+
+/**
+ *  class of managment of data on GPU global memory (allocation,desallocation & copy)
+ */
+
+class moCudaObject {
+
+public:
+
+	/**
+	 *Copy data from CPU memory to GPU global memory (default copy)
+	 *@param _data the data to allocate on GPU
+	 *@param _dataTocpy the data to copy from CPU memory to _data on GPU memory
+	 *@param _dataSize the size of data to copy
+	 */
+	template<typename T>
+	void memCopy(T* & _data, T * & _dataTocpy, unsigned _dataSize) {
+		malloc(_data, _dataSize);
+		copy(_data, _dataTocpy, _dataSize);
+	}
+
+  /**
+   *Copy device data from GPU global memory to global variable declared in device
+   *@param _dev_data the device global variable
+   *@param _dataTocpy the data to copy GPU global memory to GPU global variable
+   */
+  template<typename T>
+    void memCopyGlobalVariable(T* & _dev_data, T * & _dataTocpy) {
+    copy(_dev_data, _dataTocpy);
+  }
+
+  /**
+   *Desallocate data on GPU global memory
+   *@param _data the data to desallocate from GPU global memory
+   */
+  template<typename T>
+    void memFree(T* & _data) {
+    free(_data);
+  }
+
+public:
+
+  moCudaAllocator malloc;
+  moCudaCopy copy;
+  moCudaDesallocator free;
+
+};
+
+#endif

ParadisEO-GPU-good/src/memory/moCudaSpecificData.h

@@ -0,0 +1,85 @@
+/*
+  <moCudaSpecificData.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaSpecificData_H_
+#define _moCudaSpecificData_H_
+
+#include <memory/moCudaObject.h>
+
+/**
+ *  class of managment of specific data problem
+ */
+
+class moCudaSpecificData {
+
+ public:
+
+  /*
+   * Default constructor
+   */
+
+  moCudaSpecificData() {
+    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;
+  moCudaObject cudaObject;
+
+};
+#endif

ParadisEO-GPU-good/src/neighborhood/moCudaBitNeighbor.h

@@ -0,0 +1,121 @@
+/*
+  <moCudaBitNeighbor.h>
+  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 _moCudaBitNeighbor_h
+#define _moCudaBitNeighbor_h
+
+#include <neighborhood/moBackableNeighbor.h>
+#include <neighborhood/moIndexNeighbor.h>
+
+/**
+ * Neighbor related to a solution vector of EOT (type)
+ */
+
+template<class EOT, class Fitness = typename EOT::Fitness>
+  class moCudaBitNeighbor: public moBackableNeighbor<EOT, Fitness> ,
+  public moIndexNeighbor<EOT, Fitness> {
+
+ public:
+
+ using moBackableNeighbor<EOT, Fitness>::fitness;
+ using moIndexNeighbor<EOT, Fitness>::key;
+
+ /**
+  * move the solution
+  * @param
+  *  _solution the solution to move
+  */
+
+ virtual void move(EOT & _solution) {
+
+   _solution[key] = !_solution[key];
+
+ }
+
+ /**
+  * move back the solution (useful for the evaluation by modif)
+  * @param
+  * _solution the solution to move back
+  */
+
+ virtual void moveBack(EOT & _solution) {
+
+   _solution[key] = !_solution[key];
+
+ }
+
+ /**
+  * Return the class name.
+  * @return the class name as a std::string
+  */
+ virtual std::string className() const {
+   return "moCudaBitNeighbor";
+ }
+
+ /**
+  * 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

ParadisEO-GPU-good/src/neighborhood/moCudaKflipNeighborhood.h

@@ -0,0 +1,119 @@
+/*
+  <moKswapNeighborhood.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  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 _moCudaKswapNeighborhood_h
+#define _moCudaKswapNeighborhood_h
+
+#include <neighborhood/moKswapNeighborhood.h>
+#include <eval/moCudaEval.h>
+
+/**
+ * K-flip Neighborhood
+ */
+template<class N>
+class moCudaKflipNeighborhood: public moKswapNeighborhood<N> {
+
+ public:
+
+  /**
+   * Define a Neighbor and type of a solution corresponding
+   */
+
+  typedef N Neighbor;
+  typedef typename Neighbor::EOT EOT;
+
+  using moKswapNeighborhood<Neighbor>::neighborhoodSize;
+  using moKswapNeighborhood<Neighbor>::currentIndex;
+  using moKswapNeighborhood<Neighbor>::indices;
+  using moKswapNeighborhood<Neighbor>::mapping;
+  using moKswapNeighborhood<Neighbor>::Kswap;
+  using moKswapNeighborhood<Neighbor>::size;
+  using moKswapNeighborhood<Neighbor>::mutex;
+
+  /**
+   * Constructor
+   * @param _size the size of solution
+   * @param _size the solution size
+   * @param _Kflip the number of bit to flip
+   * @param _eval show how to evaluate neighborhood of a solution at one time
+   */
+
+ moCudaKflipNeighborhood(unsigned int _size, unsigned int _Kflip,
+			 moCudaEval<Neighbor>& _eval) :
+  moKswapNeighborhood<Neighbor> (_size, _Kflip), eval(_eval) {
+    sendMapping = true;
+    cudaMalloc((void**) &device_Mapping, sizeof(unsigned int)
+	       * neighborhoodSize * (_Kflip + 1));
+  }
+  ;
+
+  /**
+   *Destructor
+   */
+
+  ~moCudaKflipNeighborhood() {
+
+    cudaFree(device_Mapping);
+  }
+
+  /**
+   * Initialization of the neighborhood
+   * @param _solution the solution to explore
+   * @param _current the first neighbor
+   */
+
+  virtual void init(EOT& _solution, Neighbor& _current) {
+
+    moKswapNeighborhood<Neighbor>::init(_solution, _current);
+    if (sendMapping) {
+      cudaMemcpy(device_Mapping, mapping, (Kswap + 1) * neighborhoodSize
+		 * sizeof(unsigned int), cudaMemcpyHostToDevice);
+      sendMapping = false;
+    }
+    //Compute all neighbors fitness at one time
+    eval.neighborhoodKflipEval(_solution, device_Mapping, Kswap);
+  }
+
+  /**
+   * Return the class Name
+   * @return the class name as a std::string
+   */
+  virtual std::string className() const {
+    return "moCudaKSwapNeighborhood";
+  }
+
+ protected:
+
+  moCudaEval<Neighbor>& eval;
+  bool sendMapping;
+  unsigned int * device_Mapping;
+
+};
+
+#endif

ParadisEO-GPU-good/src/neighborhood/moCudaKswapNeighborhood.h

@@ -0,0 +1,119 @@
+/*
+  <moKswapNeighborhood.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Boufaras Karima, Thé Van Luong
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  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 _moCudaKswapNeighborhood_h
+#define _moCudaKswapNeighborhood_h
+
+#include <neighborhood/moKswapNeighborhood.h>
+#include <eval/moCudaEval.h>
+
+/**
+ * K-Swap Neighborhood
+ */
+template<class N>
+class moCudaKswapNeighborhood: public moKswapNeighborhood<N> {
+
+ public:
+
+  /**
+   * Define a Neighbor and type of a solution corresponding
+   */
+
+  typedef N Neighbor;
+  typedef typename Neighbor::EOT EOT;
+
+  using moKswapNeighborhood<Neighbor>::neighborhoodSize;
+  using moKswapNeighborhood<Neighbor>::currentIndex;
+  using moKswapNeighborhood<Neighbor>::indices;
+  using moKswapNeighborhood<Neighbor>::mapping;
+  using moKswapNeighborhood<Neighbor>::Kswap;
+  using moKswapNeighborhood<Neighbor>::size;
+  using moKswapNeighborhood<Neighbor>::mutex;
+
+  /**
+   * Constructor
+   * @param _size the size of solution
+   * @param _Kswap the number of swap
+   * @param _eval show how to evaluate neighborhood of a solution at one time
+   */
+
+ moCudaKswapNeighborhood(unsigned int _size, unsigned int _Kswap,
+			 moCudaEval<Neighbor>& _eval) :
+  moKswapNeighborhood<Neighbor> (_size, _Kswap), eval(_eval) {
+    sendMapping = true;
+    cudaMalloc((void**) &device_Mapping, sizeof(unsigned int)
+	       * neighborhoodSize * (Kswap + 1));
+  }
+  ;
+
+  /**
+   *Destructor
+   */
+
+  ~moCudaKswapNeighborhood() {
+
+    cudaFree(device_Mapping);
+  }
+
+  /**
+   * Initialization of the neighborhood
+   * @param _solution the solution to explore
+   * @param _current the first neighbor
+   */
+
+  virtual void init(EOT& _solution, Neighbor& _current) {
+
+    moKswapNeighborhood<Neighbor>::init(_solution, _current);
+    //Send Mapping for only the first time
+    if (sendMapping) {
+      cudaMemcpy(device_Mapping, mapping, (Kswap + 1) * neighborhoodSize
+		 * sizeof(unsigned int), cudaMemcpyHostToDevice);
+      sendMapping = false;
+    }
+    //Compute all neighbors fitness at one time
+    eval.neighborhoodKswapEval(_solution, device_Mapping, Kswap);
+  }
+
+  /**
+   * Return the class Name
+   * @return the class name as a std::string
+   */
+  virtual std::string className() const {
+    return "moCudaKSwapNeighborhood";
+  }
+
+ protected:
+
+  moCudaEval<Neighbor>& eval;
+  bool sendMapping;
+  unsigned int * device_Mapping;
+
+};
+
+#endif

ParadisEO-GPU-good/src/neighborhood/moCudaOrderNeighborhood.h

@@ -0,0 +1,96 @@
+/*
+  <moCudaOrderNeighborhood.h>
+  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 _moCudaOrderNeighborhood_h
+#define _moCudaOrderNeighborhood_h
+
+#include <neighborhood/moOrderNeighborhood.h>
+#include <eval/moCudaEval.h>
+
+/**
+ * An ordered neighborhood with parallel evaluation
+ */
+
+template<class N>
+class moCudaOrderNeighborhood: public moOrderNeighborhood<N> {
+
+ public:
+
+  /**
+   * Define type of a solution corresponding to Neighbor
+   */
+
+  typedef N Neighbor;
+  typedef typename Neighbor::EOT EOT;
+
+  using moOrderNeighborhood<Neighbor>::neighborhoodSize;
+  using moOrderNeighborhood<Neighbor>::currentIndex;
+
+  /**
+   * Constructor
+   * @param _neighborhoodSize the size of the neighborhood
+   * @param _eval show how to evaluate neighborhood of a solution at one time
+   */
+
+ moCudaOrderNeighborhood(unsigned int _neighborhoodSize,
+			 moCudaEval<Neighbor>& _eval) :
+  moOrderNeighborhood<Neighbor> (_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<Neighbor>::init(_solution, _neighbor);
+    //Compute all neighbors fitness at one time
+    eval.neighborhoodEval(_solution);
+  }
+
+  /**
+   * Return the class name.
+   * @return the class name as a std::string
+   */
+  virtual std::string className() const {
+    return "moCudaOrderNeighborhood";
+  }
+
+ protected:
+  moCudaEval<Neighbor>& eval;
+};
+
+#endif

ParadisEO-GPU-good/src/neighborhood/moCudaRndWithReplNeighborhood.h

@@ -0,0 +1,88 @@
+/*
+  <moCudaRndWithReplNeighborhood.h>
+  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 _moCudaRndWithReplNeighborhood_h
+#define _moCudaRndWithReplNeighborhood_h
+
+#include <neighborhood/moRndWithReplNeighborhood.h>
+#include <eval/moCudaEval.h>
+
+/**
+ * A Random With replacement Neighborhood with parallel evaluation
+ */
+template<class Neighbor>
+class moCudaRndWithReplNeighborhood: public moRndWithReplNeighborhood<Neighbor> {
+ public:
+
+  /**
+   * Define type of a solution corresponding to Neighbor
+   */
+  typedef typename Neighbor::EOT EOT;
+
+  using moRndWithReplNeighborhood<Neighbor>::neighborhoodSize;
+
+  /**
+   * Constructor
+   * @param _neighborhoodSize the size of the neighborhood
+   * @param _eval show how to evaluate neighborhood of a solution at one time
+   */
+ moCudaRndWithReplNeighborhood(unsigned int _neighborhoodSize, moCudaEval<
+			       Neighbor>& _eval) :
+  moRndWithReplNeighborhood<Neighbor> (_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<Neighbor>::init(_solution, _neighbor);
+    //Compute all neighbors fitness at one time
+    eval.neighborhoodEval(_solution);
+
+  }
+
+  /**
+   * Return the class Name
+   * @return the class name as a std::string
+   */
+  virtual std::string className() const {
+    return "moCudaRndWithReplNeighborhood";
+  }
+ protected:
+  moCudaEval<Neighbor>& eval;
+};
+
+#endif

ParadisEO-GPU-good/src/neighborhood/moCudaRndWithoutReplNeighborhood.h

@@ -0,0 +1,90 @@
+/*
+  <moCudaRndWithoutReplNeighborhood.h>
+  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 _moCudaRndWithoutReplNeighborhood_h
+#define _moCudaRndWithoutReplNeighborhood_h
+
+#include <neighborhood/moRndWithoutReplNeighborhood.h>
+#include <eval/moCudaEval.h>
+
+/**
+ * A Random without replacement Neighborhood with parallel evaluation
+ */
+template<class Neighbor>
+class moCudaRndWithoutReplNeighborhood: public moRndWithoutReplNeighborhood<Neighbor> {
+ public:
+
+  /**
+   * Define type of a solution corresponding to Neighbor
+   */
+  typedef typename Neighbor::EOT EOT;
+
+  using moRndWithoutReplNeighborhood<Neighbor>::neighborhoodSize;
+  using moRndWithoutReplNeighborhood<Neighbor>::maxIndex;
+  using moRndWithoutReplNeighborhood<Neighbor>::indexVector;
+  /**
+   * Constructor
+   * @param _neighborhoodSize the size of the neighborhood
+   * @param _eval show how to evaluate neighborhood of a solution at one time
+   */
+ moCudaRndWithoutReplNeighborhood(unsigned int _neighborhoodSize,moCudaEval<
+				  Neighbor>& _eval) :
+  moRndWithoutReplNeighborhood<Neighbor> (_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<Neighbor>::init(_solution, _neighbor);
+    //Compute all neighbors fitness at one time
+    eval.neighborhoodEval(_solution);
+  }
+
+  /**
+   * Return the class Name
+   * @return the class name as a std::string
+   */
+  virtual std::string className() const {
+    return "moCudaRndWithoutReplNeighborhood";
+  }
+
+ protected:
+  moCudaEval<Neighbor>& eval;
+};
+
+#endif

ParadisEO-GPU-good/src/performance/moCudaTimer.h

@@ -0,0 +1,116 @@
+/*
+  <moCudaTimer.h>
+  Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
+
+  Karima Boufaras, Thé Van LUONG
+
+  This software is governed by the CeCILL license under French law and
+  abiding by the rules of distribution of free software.  You can  use,
+  modify and/ or redistribute the software under the terms of the CeCILL
+  license as circulated by CEA, CNRS and INRIA at the following URL
+  "http://www.cecill.info".
+
+  As a counterpart to the access to the source code and  rights to copy,
+  modify and redistribute granted by the license, users are provided only
+  with a limited warranty  and the software's author,  the holder of the
+  economic rights,  and the successive licensors  have only  limited liability.
+
+  In this respect, the user's attention is drawn to the risks associated
+  with loading,  using,  modifying and/or developing or reproducing the
+  software by the user in light of its specific status of free software,
+  that may mean  that it is complicated to manipulate,  and  that  also
+  therefore means  that it is reserved for developers  and  experienced
+  professionals having in-depth computer knowledge. Users are therefore
+  encouraged to load and test the software's suitability as regards their
+  requirements in conditions enabling the security of their systems and/or
+  data to be ensured and,  more generally, to use and operate it in the
+  same conditions as regards security.
+  The fact that you are presently reading this means that you have had
+  knowledge of the CeCILL license and that you accept its terms.
+
+  ParadisEO WebSite : http://paradiseo.gforge.inria.fr
+  Contact: paradiseo-help@lists.gforge.inria.fr
+*/
+
+#ifndef _moCudaTimer_H_
+#define _moCudaTimer_H_
+
+// CUDA Utility Tools 
+#include <cutil.h>
+
+/**
+ * To compute execution time
+ */
+
+class moCudaTimer {
+
+ public:
+
+  unsigned int timer;
+
+  /**
+   * Constructor
+   */
+
+  moCudaTimer() {
+
+    timer = 0;
+    //Create timer
+    cutCreateTimer(&timer);
+  }
+
+  /**
+   * Destructor
+   */
+
+  ~moCudaTimer() {
+
+  }
+
+  /**
+   * to start compute execution time
+   */
+
+  void start() {
+
+    // Start timer
+    cutStartTimer(timer);
+  }
+
+  /**
+   * to stop compute execution time
+   */
+
+  void stop() {
+
+    // Stop timer
+    cutStopTimer(timer);
+
+  }
+
+  /**
+   * to get timer value in ms
+   * @return execution time in ms
+   */
+
+  double getTime() {
+
+    // get time
+    return cutGetTimerValue(timer);
+
+  }
+
+  /**
+   * to delete timer
+   */
+
+  void deleteTimer() {
+
+    //Delete timer
+    cutDeleteTimer(timer);
+
+  }
+
+};
+
+#endif