From 0ab74e3bb2fed53e2a8000f73ab997669c366464 Mon Sep 17 00:00:00 2001
From: boufaras <boufaras@331e1502-861f-0410-8da2-ba01fb791d7f>
Date: Fri, 28 Jan 2011 09:51:56 +0000
Subject: [PATCH] git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2083
 331e1502-861f-0410-8da2-ba01fb791d7f

---
 ParadisEO-GPU/src/eval/moCudaKswapEval.h | 21 ++++++++++++---------
 1 file changed, 12 insertions(+), 9 deletions(-)

diff --git a/ParadisEO-GPU/src/eval/moCudaKswapEval.h b/ParadisEO-GPU/src/eval/moCudaKswapEval.h
index ea7fe3990..084d9e73d 100644
--- a/ParadisEO-GPU/src/eval/moCudaKswapEval.h
+++ b/ParadisEO-GPU/src/eval/moCudaKswapEval.h
@@ -38,7 +38,7 @@
 #include <eval/moCudakernelEval.h>
 
 /**
- * class for the cuda evaluation
+ * class for the K-swap neighborhood evaluation
  */
 
 template<class Neighbor, class IncrementEval>
@@ -75,7 +75,6 @@ public:
 		moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
 		mutex = false;
 		mutex_kswap=false;
-		compt=0;
 	}
 	/**
 	 * Destructor
@@ -99,7 +98,7 @@ public:
 	/**
 	 * Compute fitness for all solution neighbors in device
 	 * @param _sol the solution which generate the neighborhood
-	 * @param _mapping the array of indices mapping
+	 * @param _mapping the array of mapping indexes for K-swap neighborhood
 	 * @param _Kswap the number of swap
 	 */
 
@@ -110,6 +109,7 @@ public:
 
 		// Get Current solution fitness
 		Fitness fitness = _sol.fitness();
+
 		//Case of Permutation
 		if (_Kswap == 1) {
 			if (!mutex) {
@@ -117,6 +117,7 @@ public:
 				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);
@@ -130,6 +131,7 @@ public:
 		//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
@@ -138,6 +140,7 @@ public:
 				cudaMalloc((void**) &device_tmp.vect, neighborhoodSize
 						* sizeof(T));
 				mutex_kswap = true;
+
 			}
 
 			for (int i = 0; i < neighborhoodSize; i++) {
@@ -145,19 +148,17 @@ public:
 					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);
-		   std::cout<<" Iteration number "<<compt<<" "<<std::endl;
-		/*	for(unsigned int i=0;i<neighborhoodSize;i++)
-				std::cout<<" "<<host_FitnessArray[i]<<" ";//<<std::endl;*/
-			compt++;
 		}
 
 	}
@@ -165,7 +166,7 @@ public:
 	/**
 	 * 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 indices mapping
+	 * @param _mapping the array of mapping indexes for k-flip neighborhood
 	 * @param _Kflip the number of flip to do
 	 */
 
@@ -176,17 +177,20 @@ public:
 
 		// 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);
@@ -203,7 +207,6 @@ protected:
 	T * vect;
 	bool mutex_kswap;
 	bool mutex;
-	unsigned int compt;
 };
 
 #endif