added missing sync for async copies

Makefile

@@ -205,8 +205,8 @@ PDB      := 3ce3
 NRUN     := 100
 POPSIZE  := 150
 TESTNAME := test
-TESTLS   := sd
-NUM_LSIT := 30
+TESTLS   := sw
+NUM_LSIT := 300
 
 test: odock
 	$(BIN_DIR)/$(TARGET) -ffile ./input/$(PDB)/derived/$(PDB)_protein.maps.fld -lfile ./input/$(PDB)/derived/$(PDB)_ligand.pdbqt -nrun $(NRUN) -psize $(POPSIZE) -resnam $(TESTNAME) -gfpop 1 -lsmet $(TESTLS) -lsit $(NUM_LSIT) -smooth 0.5

device/kernel1.cl

@@ -81,13 +81,14 @@ gpu_calc_initpop(
 			                   dockpars_conformations_current + GENOTYPE_LENGTH_IN_GLOBMEM*get_group_id(0),
 			                   GENOTYPE_LENGTH_IN_GLOBMEM, 0);
 
-	wait_group_events(1,&ev);
-
 	// Determining run-ID
 	if (get_local_id(0) == 0) {
 		run_id = get_group_id(0) / dockpars_pop_size;
 	}
 
+	// Asynchronous copy should be finished by here
+	wait_group_events(1,&ev);
+
 	// Evaluating initial genotypes
 	barrier(CLK_LOCAL_MEM_FENCE);
 

device/kernel2.cl

@@ -43,10 +43,12 @@ gpu_sum_evals(
 	// Maximum size defined in ../common/defines.h
 	__local int local_evals_of_new_entities[MAX_POPSIZE];
 
-	async_work_group_copy(local_evals_of_new_entities,
-                              dockpars_evals_of_new_entities+get_group_id(0)*pop_size,
-                              pop_size, 0);
+	event_t ev =async_work_group_copy(local_evals_of_new_entities,
+                              		  dockpars_evals_of_new_entities+get_group_id(0)*pop_size,
+                                          pop_size, 0);
 
+	// Asynchronous copy should be finished by here
+	wait_group_events(1,&ev);
 
   	for (entity_counter = get_local_id(0);
 	     entity_counter < pop_size;

device/kernel3.cl

@@ -130,9 +130,9 @@ perform_LS(
 
 	barrier(CLK_LOCAL_MEM_FENCE);
 
-  	async_work_group_copy(offspring_genotype,
-			      dockpars_conformations_next+(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM,
-                              dockpars_num_of_genes, 0);
+  	event_t ev = async_work_group_copy(offspring_genotype,
+			      		   dockpars_conformations_next+(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM,
+                              		   dockpars_num_of_genes, 0);
 
 	for (gene_counter = get_local_id(0);
 	     gene_counter < dockpars_num_of_genes;
@@ -148,6 +148,10 @@ perform_LS(
 		evaluation_cnt = 0;
 	}
 
+
+	// Asynchronous copy should be finished by here
+	wait_group_events(1,&ev);
+
 	barrier(CLK_LOCAL_MEM_FENCE);
 
 	while ((iteration_cnt < dockpars_max_num_of_iters) && (rho > dockpars_rho_lower_bound))

device/kernel4.cl

@@ -211,9 +211,12 @@ gpu_gen_and_eval_newpops(
 		}
 		else	//no crossover
 		{
-			async_work_group_copy(offspring_genotype,
-					     dockpars_conformations_current+(run_id*dockpars_pop_size+parents[0])*GENOTYPE_LENGTH_IN_GLOBMEM,
-					     dockpars_num_of_genes, 0);
+			event_t ev = async_work_group_copy(offspring_genotype,
+					     		   dockpars_conformations_current+(run_id*dockpars_pop_size+parents[0])*GENOTYPE_LENGTH_IN_GLOBMEM,
+					                   dockpars_num_of_genes, 0);
+
+			// Asynchronous copy should be finished by here
+			wait_group_events(1,&ev);
 		} // End of crossover
 
 		barrier(CLK_LOCAL_MEM_FENCE);

device/kernel_gradient.cl

@@ -126,9 +126,9 @@ gradient_minimizer(
 
 	barrier(CLK_LOCAL_MEM_FENCE);
 
-  	async_work_group_copy(genotype,
-  			      dockpars_conformations_next+(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM,
-                              dockpars_num_of_genes, 0);
+  	event_t ev = async_work_group_copy(genotype,
+  			      		   dockpars_conformations_next+(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM,
+                              		   dockpars_num_of_genes, 0);
 
   	// -----------------------------------------------------------------------------
   	// Some OpenCL compilers don't allow declaring 
@@ -218,7 +218,6 @@ gradient_minimizer(
 	float ligand_center_z = -2.40896296;
 	#endif
 
-
 	// Defining lower and upper bounds for genotypes
 	__local float lower_bounds_genotype[ACTUAL_GENOTYPE_LENGTH];
 	__local float upper_bounds_genotype[ACTUAL_GENOTYPE_LENGTH];
@@ -262,6 +261,9 @@ gradient_minimizer(
 	__local float torsions_genotype[ACTUAL_GENOTYPE_LENGTH];
 	//----------------------------------
 
+	// Asynchronous copy should be finished by here
+	wait_group_events(1,&ev);
+
 	// The termination criteria is based on 
 	// a maximum number of iterations, and
 	// the minimum step size allowed for single-floating point numbers