more fixes; get coverage quickly

packages/cuda/cuda.c

@@ -268,6 +268,48 @@ cuda_eval( void )
   return YAP_Unify(YAP_ARG2, out);
 }
 
+static int
+cuda_coverage( void )
+{
+  int32_t *mat;
+  predicate *ptr = (predicate *)YAP_IntOfTerm(YAP_ARG1);
+  int32_t n = Cuda_Eval(facts, cf, rules, cr, ptr, & mat);
+  int32_t ncols = ptr->num_columns;
+  int32_t post = YAP_AtomToInt(YAP_AtomOfTerm(YAP_ARG2));
+  int32_t i = n/2, min = 0, max = n-1, t0 = mat[0], t1 = mat[(n-1)*2];
+
+  if (n < 0)
+    return FALSE;
+  if (t0 == t1) { /* all sametype */
+    free( mat );
+    /* all pos */
+    if (t0 == post) 
+      return YAP_Unify(YAP_ARG3, YAP_MkIntTerm(n)) && 
+	YAP_Unify(YAP_ARG4, YAP_MkIntTerm(0));
+    /* all neg */
+    return YAP_Unify(YAP_ARG4, YAP_MkIntTerm(n)) && 
+      YAP_Unify(YAP_ARG3, YAP_MkIntTerm(0));
+  }
+  do {
+    i = (min+max)/2;
+    if (i == min) i++;
+    if (mat[i*2] == t0) {
+      min = i;
+    } else {
+      max = i;
+    }
+    if (min+1 == max) {      
+      free( mat );
+      if (t0 == post) 
+	return YAP_Unify(YAP_ARG3, YAP_MkIntTerm(max)) && 
+	  YAP_Unify(YAP_ARG4, YAP_MkIntTerm(n-max));
+      /* all neg */
+      return YAP_Unify(YAP_ARG4, YAP_MkIntTerm(max)) && 
+	YAP_Unify(YAP_ARG3, YAP_MkIntTerm(n-max));
+    }
+  } while ( TRUE );
+}
+
 static int cuda_count( void )
 {
   int32_t *mat;
@@ -298,6 +340,7 @@ init_cuda(void)
   YAP_UserCPredicate("load_rule", load_rule, 4);
   YAP_UserCPredicate("cuda_erase", cuda_erase, 1);
   YAP_UserCPredicate("cuda_eval", cuda_eval, 2);
+  YAP_UserCPredicate("cuda_coverage", cuda_coverage, 4);
   YAP_UserCPredicate("cuda_count", cuda_count, 2);
 }
 

packages/cuda/cuda.yap

@@ -3,6 +3,7 @@
 		 cuda_rule/2,
 		 cuda_erase/1,
 		 cuda_eval/2,
+		 cuda_coverage/4,
 		 cuda_count/2]).
 
 tell_warning :-

packages/cuda/lista.cu

@@ -67,6 +67,73 @@ void buscarreglas(vector<gpunode> *facts, vector<rulenode> *rules)
 	}
 }
 
+template<class InputIterator>
+void movebpreds(InputIterator rules, InputIterator end)
+{
+	int x, subs, total, cont, cont2, pos;
+	int *move, *rest;
+	while(rules != end)
+	{
+		if(rules->num_bpreds.x > 0)
+		{
+
+			total = rules->num_rows+rules->num_bpreds.x;
+
+			/*cout << "ANTES" << endl;
+			for(x = 0; x < rules->rule_names[total]; x++)
+				cout << rules->address_host_table[x] << " ";
+			cout << "FINANTES" << endl;*/
+
+			move = (int *)malloc(sizeof(int) * rules->num_bpreds.x * 4);
+			rest = (int *)malloc(sizeof(int) * rules->rule_names[total]);
+			cont = 0;
+			cont2 = 0;
+			for(x = 0; x < total; x++)
+			{
+				subs = rules->rule_names[x+1] - rules->rule_names[x];
+				
+				//cout << subs << " ";
+				
+				if(rules->address_host_table[rules->rule_names[x]] > 0)
+				{
+					memcpy(rest + cont, rules->address_host_table + rules->rule_names[x], subs * sizeof(int));
+					cont += subs;
+				}
+				else
+				{
+					memcpy(move + cont2, rules->address_host_table + rules->rule_names[x], subs * sizeof(int));
+					cont2 += subs;
+				}
+			}
+
+			/*cout << "REST" << endl;
+			for(x = 0; x < cont; x++)
+				cout << rest[x] << " ";
+			cout << "RESTFIN" << endl;*/
+
+			memcpy(rest + cont, move, cont2 * sizeof(int));
+			pos = 1;
+			for(x = 1; x <= total; x++)
+			{
+				while(rest[pos] != 0)
+					pos++;
+				pos++;
+				rules->rule_names[x] = pos;
+			}
+			memcpy(rules->address_host_table, rest, sizeof(int) * rules->rule_names[total]);
+			free(move);
+			free(rest);
+
+			/*cout << "DESPUES" << endl;
+			for(x = 0; x < rules->rule_names[total]; x++)
+				cout << rules->address_host_table[x] << " ";
+			cout << "FINDESPUES" << endl;*/
+
+		}
+		rules++;
+	}
+}
+
 template<class InputIterator>
 void nombres(InputIterator rules, InputIterator end)
 {
@@ -454,7 +521,7 @@ void proyeccion(InputIterator actual, InputIterator end)
 			pos = columnsproject(pv, pos.y, actual->address_host_table, ini, fin, rulestart, ruleend, &res, &pv);
 			actual->project[numjoins] = res;
 			actual->projpos[numjoins] = pos;
-			actual->num_bpreds.y = pos.y; /*para guardar el tamaño de la union final*/
+			actual->num_bpreds.y = pos.y; /*para guardar el tamanio de la union final*/
 			actual->num_bpreds.z = builtinpredicates(pv, pos.y, actual->address_host_table, ruleend + 1, actual->rule_names[total] - 1, &res);
 			actual->builtin = res;
 		}
@@ -797,6 +864,11 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 	for(x = 0; x < ninpr; x++)
 		L.push_back(*inprules[x]);
 
+	/*cout << "NAMES" << endl;
+	for(x = 0; x < (ninpf+ninpr); x++)
+		cout << L[x].name << endl;
+	cout << "NAMESEND" << endl;*/
+
 	qname = inpquery->name;
 	query = inpquery->address_host_table;
 	qsize = inpquery->num_columns;
@@ -822,7 +894,8 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 	rul_str = rules.begin();
 	fin = rules.end();
 
-	nombres(rul_str, fin);
+	nombres(rul_str, fin); /*preprocessing*/
+	movebpreds(rul_str, fin);
 	referencias(L.begin(), L.end(), rul_str, fin);
 	seleccion(rul_str, fin);
 	selfjoin(rul_str, fin);
@@ -1027,6 +1100,8 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 					cudaEventRecord(stop3, 0);
 					cudaEventSynchronize(stop3);
 					cudaEventElapsedTime(&time, start3, stop3);
+					cudaEventDestroy(start3);
+					cudaEventDestroy(stop3);
 					cout << "Predicados = " << time << endl;
 					#endif
 				}
@@ -1046,6 +1121,8 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 				cudaEventRecord(stop2, 0);
 				cudaEventSynchronize(stop2);
 				cudaEventElapsedTime(&time, start2, stop2);
+				cudaEventDestroy(start2);
+				cudaEventDestroy(stop2);
 				cout << "Union = " << time << endl;
 				#endif					
 	
@@ -1175,6 +1252,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 		tipo = res_rows * cols1 * sizeof(int);
 		hres = (int *)malloc(tipo);
 		cudaMemcpy(hres, res, tipo, cudaMemcpyDeviceToHost);
+		cudaFree(res);
 	}
 	else
 		res_rows = 0;
@@ -1182,6 +1260,8 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 	cudaEventRecord(stop, 0);
 	cudaEventSynchronize(stop);
 	cudaEventElapsedTime(&time, start, stop);
+	cudaEventDestroy(start);
+	cudaEventDestroy(stop);
 
 	if(showr == 1)
 	{
@@ -1199,6 +1279,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr,
 	cout << "Size = " << res_rows << endl;
 	cout << "Iterations = " << itr << endl;
 
+	clear_memory();
 	*result = hres;
 
 	return res_rows;

packages/cuda/memory.cu

@@ -182,7 +182,7 @@ void liberar(int *ptr, int size)
 
 void reservar(int **ptr, int size)
 {
-	//cout << "R " << avmem << " " << size;
+        // cout << "R " << avmem << " " << size
 
 	while(avmem < size)
 		limpiar();
@@ -190,7 +190,7 @@ void reservar(int **ptr, int size)
 		limpiar();
 	avmem -= size;
 
-	//cout << " " << avmem << endl;
+	// cout << " " << avmem << endl;
 }
 
 void registrar(int name, int num_columns, int *ptr, int rows, int itr, int rule)
@@ -346,9 +346,6 @@ int cargafinal(int name, int cols, int **ptr)
 	}
 	cout << "select finala" << endl;*/
 
-	GPUmem.clear();
-	CPUmem.clear();
-
 	*ptr = ini;
 	return cont;
 }
@@ -463,22 +460,24 @@ void resultados(vector<rulenode>::iterator first, vector<rulenode>::iterator las
 	cout << cont << endl;
 }
 
-/*device_vector<int> reservar_vector(int size)
+void clear_memory()
 {
-	limpiar(size * sizeof(int));
+	list<memnode>::iterator ini;
-	device_vector<int> ret(size);
+	list<memnode>::iterator fin;
-	return ret;
+	ini = GPUmem.begin();
-}*/
+	fin = GPUmem.end();
-/*
+	while(ini != fin)
-void reservar_resultado(InputIterator req, int *ptr, int size)
+	{
-{
+		cudaFree(ini->dev_address);
-	limpiar(size);
+		ini++;
-	memnode temp;
+	}
-	temp.name = req->name;
+	ini = CPUmem.begin();
-	temp.size = size;
+	fin = CPUmem.end();
-	cudaMalloc(&temp.dev_address, size);
+	while(ini != fin)
-	temp.in_use = 1;
+	{
-	GPUmem.push_back(temp);
+		free(ini->dev_address);
-	avmem -= size;
+		ini++;
-	ptr = temp.dev_address;
+	}
-}*/
+	GPUmem.clear();
+	CPUmem.clear();
+}