From 2250ff6e11ca18c79fd734ea3ab4204b8a4b3114 Mon Sep 17 00:00:00 2001
From: boufaras <boufaras@331e1502-861f-0410-8da2-ba01fb791d7f>
Date: Tue, 11 Jan 2011 16:05:02 +0000
Subject: [PATCH] git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2062
 331e1502-861f-0410-8da2-ba01fb791d7f

---
 ParadisEO-GPU/src/eval/moCudaKswapEval.h | 68 +++++++++++++++++-------
 1 file changed, 49 insertions(+), 19 deletions(-)

diff --git a/ParadisEO-GPU/src/eval/moCudaKswapEval.h b/ParadisEO-GPU/src/eval/moCudaKswapEval.h
index 75ae3acb8..ea7fe3990 100644
--- a/ParadisEO-GPU/src/eval/moCudaKswapEval.h
+++ b/ParadisEO-GPU/src/eval/moCudaKswapEval.h
@@ -73,6 +73,19 @@ public:
 
 	moCudaKswapEval(unsigned int _neighborhoodSize, IncrementEval & _incrEval) :
 		moCudaEval<Neighbor> (_neighborhoodSize), incrEval(_incrEval) {
+		mutex = false;
+		mutex_kswap=false;
+		compt=0;
+	}
+	/**
+	 * Destructor
+	 */
+	~moCudaKswapEval() {
+		if(mutex_kswap){
+		cudaFree(&device_setSolution);
+		cudaFree(&device_tmp);
+		delete[] vect;
+		}
 	}
 
 	/**
@@ -99,8 +112,11 @@ public:
 		Fitness fitness = _sol.fitness();
 		//Case of Permutation
 		if (_Kswap == 1) {
-			//Allocate the space for solution in the device global memory
-			cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
+			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);
@@ -113,33 +129,37 @@ public:
 		}
 		//Case Kswap
 		else if (_Kswap > 1) {
-			//NeighborhoodSize copy of solution
-			EOT device_setSolution;
-			//NeighborhoodSize element of EOT
-			EOT device_tmp;
-			//Vector of neighborhoodSize copy of solution
-			T * vect = new T[neighborhoodSize * _size];
+			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];
 				}
 			}
-			//Allocate the space for set of solution in the device global memory
-			cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize
-					* _size * sizeof(T));
-
 			//Copy the set of solution from the host to device
 			cudaMemcpy(device_setSolution.vect, vect, neighborhoodSize * _size
 					* sizeof(T), cudaMemcpyHostToDevice);
-			//Allocate the space to save temporary EOT element to swap
-			cudaMalloc((void**) &device_tmp.vect, neighborhoodSize * sizeof(T));
+
 			//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++;
 		}
+
 	}
 
 	/**
@@ -156,10 +176,11 @@ public:
 
 		// Get Current solution fitness
 		Fitness fitness = _sol.fitness();
-
-		//Allocate the space for solution in the device global memory
-		cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
-
+		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);
@@ -174,6 +195,15 @@ public:
 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;
+	unsigned int compt;
 };
 
 #endif