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

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
+			if (!mutex) {
-			cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
+				//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
+			if (!mutex_kswap) {
-			EOT device_setSolution;
+				vect = new T[neighborhoodSize * _size];
-			//NeighborhoodSize element of EOT
+				//Allocate the space for set of solution in the device global memory
-			EOT device_tmp;
+				cudaMalloc((void**) &device_setSolution.vect, neighborhoodSize
-			//Vector of neighborhoodSize copy of solution
+						* _size * sizeof(T));
-			T * vect = new T[neighborhoodSize * _size];
+				//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();
-
+		if (!mutex) {
-		//Allocate the space for solution in the device global memory
+			//Allocate the space for solution in the device global memory
-		cudaMalloc((void**) &device_solution.vect, _size * sizeof(T));
+			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