diff --git a/C/errors.c b/C/errors.c index a01ed13be..46839fffa 100644 --- a/C/errors.c +++ b/C/errors.c @@ -40,10 +40,9 @@ Yap_RestartYap ( int flag ) siglongjmp(LOCAL_RestartEnv,1); } -#ifdef DEBUG +void DumpActiveGoals( CACHE_TYPE1 ); static int hidden(Atom); static int legal_env(CELL * CACHE_TYPE); -void DumpActiveGoals( CACHE_TYPE1 ); static void detect_bug_location(yamop *,find_pred_type,char *, int); #define ONHEAP(ptr) (CellPtr(ptr) >= CellPtr(Yap_HeapBase) && CellPtr(ptr) < CellPtr(HeapTop)) @@ -86,6 +85,19 @@ legal_env (CELL *ep USES_REGS) return (TRUE); } +static int +YapPutc(int sno, wchar_t ch) +{ + return (putc(ch, GLOBAL_stderr)); +} + +static void +YapPlWrite(Term t) +{ + Yap_plwrite(t, NULL, 15, 0, 1200); +} + + void DumpActiveGoals ( USES_REGS1 ) { @@ -127,16 +139,16 @@ DumpActiveGoals ( USES_REGS1 ) if (first++ == 1) fprintf(stderr,"Active ancestors:\n"); if (pe->ModuleOfPred) mod = pe->ModuleOfPred; - Yap_DebugPlWrite (mod); - Yap_DebugPutc (LOCAL_c_error_stream,':'); + YapPlWrite (mod); + YapPutc (LOCAL_c_error_stream,':'); if (pe->ArityOfPE == 0) { - Yap_DebugPlWrite (MkAtomTerm ((Atom)f)); + YapPlWrite (MkAtomTerm ((Atom)f)); } else { - Yap_DebugPlWrite (MkAtomTerm (NameOfFunctor (f))); - Yap_DebugPutc (LOCAL_c_error_stream,'/'); - Yap_DebugPlWrite (MkIntTerm (ArityOfFunctor (f))); + YapPlWrite (MkAtomTerm (NameOfFunctor (f))); + YapPutc (LOCAL_c_error_stream,'/'); + YapPlWrite (MkIntTerm (ArityOfFunctor (f))); } - Yap_DebugPutc (LOCAL_c_error_stream,'\n'); + YapPutc (LOCAL_c_error_stream,'\n'); } else { UNLOCK(pe->PELock); } @@ -161,28 +173,28 @@ DumpActiveGoals ( USES_REGS1 ) if (pe->ModuleOfPred) mod = pe->ModuleOfPred; else mod = TermProlog; - Yap_DebugPlWrite (mod); - Yap_DebugPutc (LOCAL_c_error_stream,':'); + YapPlWrite (mod); + YapPutc (LOCAL_c_error_stream,':'); if (pe->ArityOfPE == 0) { - Yap_DebugPlWrite (MkAtomTerm (NameOfFunctor(f))); + YapPlWrite (MkAtomTerm (NameOfFunctor(f))); } else { Int i = 0, arity = pe->ArityOfPE; Term *args = &(b_ptr->cp_a1); - Yap_DebugPlWrite (MkAtomTerm (NameOfFunctor (f))); - Yap_DebugPutc (LOCAL_c_error_stream,'('); + YapPlWrite (MkAtomTerm (NameOfFunctor (f))); + YapPutc (LOCAL_c_error_stream,'('); for (i= 0; i < arity; i++) { - if (i > 0) Yap_DebugPutc (LOCAL_c_error_stream,','); - Yap_DebugPlWrite(args[i]); + if (i > 0) YapPutc (LOCAL_c_error_stream,','); + YapPlWrite(args[i]); } - Yap_DebugPutc (LOCAL_c_error_stream,')'); + YapPutc (LOCAL_c_error_stream,')'); } - Yap_DebugPutc (LOCAL_c_error_stream,'\n'); + YapPutc (LOCAL_c_error_stream,'\n'); } UNLOCK(pe->PELock); b_ptr = b_ptr->cp_b; } } -#endif /* DEBUG */ + static void detect_bug_location(yamop *yap_pc, find_pred_type where_from, char *tp, int psize) @@ -477,6 +489,11 @@ Yap_Error(yap_error_number type, Term where, char *format,...) fprintf(stderr,"%% YAP OOOPS: %s.\n",tmpbuf); fprintf(stderr,"%%\n%%\n"); } + detect_bug_location(P, FIND_PRED_FROM_ANYWHERE, (char *)H, 256); + fprintf (stderr,"%%\n%% PC: %s\n",(char *)H); + detect_bug_location(CP, FIND_PRED_FROM_ANYWHERE, (char *)H, 256); + fprintf (stderr,"%% Continuation: %s\n",(char *)H); + DumpActiveGoals( PASS_REGS1 ); error_exit_yap (1); } if (P == (yamop *)(FAILCODE)) diff --git a/configure b/configure index 5c0e235d5..4323e75bd 100755 --- a/configure +++ b/configure @@ -6969,12 +6969,12 @@ fi then ENABLE_CUDA="@# " else - ENABLE_CUDA="@# " + ENABLE_CUDA="" case "$target_os" in *darwin*) CUDA_LDFLAGS="$LDFLAGS" CUDA_CPPFLAGS="-shared -arch=sm_20 -Xcompiler -fPIC -O3 " - CUDA_SHLIB_LD="$NVCC -Xcompiler -dynamiclib -L../.. -lYAP " + CUDA_SHLIB_LD="$NVCC -Xcompiler -dynamiclib -L../.. -lYap " ;; **) CUDA_LDFLAGS="$LDFLAGS $LIBS" diff --git a/configure.in b/configure.in index dcd5fe5c1..1c2650a39 100755 --- a/configure.in +++ b/configure.in @@ -915,7 +915,7 @@ else *darwin*) CUDA_LDFLAGS="$LDFLAGS" CUDA_CPPFLAGS="-shared -arch=sm_20 -Xcompiler -fPIC -O3 " - CUDA_SHLIB_LD="$NVCC -Xcompiler -dynamiclib -L../.. -lYAP " + CUDA_SHLIB_LD="$NVCC -Xcompiler -dynamiclib -L../.. -lYap " ;; **) CUDA_LDFLAGS="$LDFLAGS $LIBS" diff --git a/library/dialect/swi/fli/swi.h b/library/dialect/swi/fli/swi.h index f47e6b221..58142e979 100644 --- a/library/dialect/swi/fli/swi.h +++ b/library/dialect/swi/fli/swi.h @@ -85,7 +85,6 @@ in_hash(ADDR key) static inline atom_t AtomToSWIAtom(Atom at) { - atom_t ats; TranslationEntry *p; if ((p = Yap_GetTranslationProp(at)) != NULL) @@ -96,7 +95,7 @@ AtomToSWIAtom(Atom at) static inline Atom SWIAtomToAtom(atom_t at) { - if ((CELL)at < 2*N_SWI_ATOMS+1) + if ((CELL)at & 1) return SWI_Atoms[at/2]; return (Atom)at; } diff --git a/packages/cuda/bpreds.cu b/packages/cuda/bpreds.cu index ea94b5b2a..1d8deadd1 100644 --- a/packages/cuda/bpreds.cu +++ b/packages/cuda/bpreds.cu @@ -46,7 +46,7 @@ __global__ void predicates(int *dop1, int rows, int cols, int *cons, int numc, i } } -int bpreds(int *dop1, int rows, int *bin, int3 numpreds, int **ret) +int bpreds(int *dop1, int rows, int cols, int *bin, int3 numpreds, int **ret) { int *temp; int tmplen = rows + 1; @@ -124,6 +124,8 @@ int bpreds(int *dop1, int rows, int *bin, int3 numpreds, int **ret) liberar(dhead, hsize); liberar(temp, size); + liberar(dop1, rows * cols * sizeof(int)); + *ret = fres; return num; } diff --git a/packages/cuda/cuda.c b/packages/cuda/cuda.c index 0010b7eac..aa49386c5 100644 --- a/packages/cuda/cuda.c +++ b/packages/cuda/cuda.c @@ -158,6 +158,69 @@ load_facts( void ) { } } +static int currentFact = 0; +static predicate *currentPred = NULL; + +static int +cuda_init_facts( void ) { + + int32_t nrows = YAP_IntOfTerm(YAP_ARG1); + int32_t ncols = YAP_IntOfTerm(YAP_ARG2), i = 0; + int32_t *mat = (int32_t *)malloc(sizeof(int32_t)*nrows*ncols); + int32_t pname = YAP_AtomToInt(YAP_AtomOfTerm(YAP_ARG3)); + predicate *pred; + + if (!mat) + return FALSE; + if (YAP_IsVarTerm( YAP_ARG4)) { + // new + pred = (predicate *)malloc(sizeof(predicate)); + } else { + pred = (predicate *)YAP_IntOfTerm(YAP_ARG4); + if (pred->address_host_table) + free( pred->address_host_table ); +} + pred->name = pname; + pred->num_rows = nrows; + pred->num_columns = ncols; + pred->is_fact = TRUE; + pred->address_host_table = mat; + currentPred = pred; + currentFact = 0; + + if (YAP_IsVarTerm( YAP_ARG4)) { + return YAP_Unify(YAP_ARG4, YAP_MkIntTerm((YAP_Int)pred)); + } else { + return TRUE; + } +} + +static int +cuda_load_fact( void ) { + YAP_Term th = YAP_ARG1; + + int i, j; + int ncols = currentPred->num_columns; + int *mat = currentPred->address_host_table; + i = currentFact; + for (j = 0; j < ncols; j++) { + YAP_Term ta = YAP_ArgOfTerm(j+1, th); + if (YAP_IsAtomTerm(ta)) { + mat[i*ncols+j] = YAP_AtomToInt(YAP_AtomOfTerm(ta)); + } else { + mat[i*ncols+j] = YAP_IntOfTerm(ta); + } + } + i++; + if (i == currentPred->num_rows) { + Cuda_NewFacts(currentPred); + currentPred = NULL; + currentFact = 0; + } else { + currentFact = i; + } +} + static int load_rule( void ) { // maximum of 2K symbols per rule, should be enough for ILP @@ -264,7 +327,8 @@ cuda_eval( void ) } out = YAP_MkPairTerm(YAP_MkApplTerm( f, ncols, vec ), out); } - free( mat ); + if (n > 0) + free( mat ); return YAP_Unify(YAP_ARG2, out); } @@ -276,10 +340,16 @@ cuda_coverage( void ) 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]; + int32_t i = n/2, min = 0, max = n-1; + int32_t t0, t1; if (n < 0) return FALSE; + if (n == 0) { + return YAP_Unify(YAP_ARG4, YAP_MkIntTerm(0)) && + YAP_Unify(YAP_ARG3, YAP_MkIntTerm(0)); + } + t0 = mat[0], t1 = mat[(n-1)*2]; if (t0 == t1) { /* all sametype */ free( mat ); /* all pos */ @@ -337,6 +407,8 @@ init_cuda(void) AtomLe = YAP_LookupAtom("=<"); AtomDf = YAP_LookupAtom("\\="); YAP_UserCPredicate("load_facts", load_facts, 4); + YAP_UserCPredicate("cuda_init_facts", cuda_init_facts, 4); + YAP_UserCPredicate("cuda_load_fact", cuda_load_fact, 1); YAP_UserCPredicate("load_rule", load_rule, 4); YAP_UserCPredicate("cuda_erase", cuda_erase, 1); YAP_UserCPredicate("cuda_eval", cuda_eval, 2); diff --git a/packages/cuda/cuda.yap b/packages/cuda/cuda.yap index 52dd03c3d..4649fc5d7 100644 --- a/packages/cuda/cuda.yap +++ b/packages/cuda/cuda.yap @@ -21,9 +21,21 @@ cuda_inline(P, Q) :- cuda_extensional( Call, IdFacts) :- strip_module(Call, Mod, Name/Arity), functor(S, Name, Arity), - findall( S, Mod:S, L), - length( L, N ), - load_facts( N, Arity, L, IdFacts ). + count_answers( Mod:S, N), + % reserve space + cuda_init_facts( N, Arity, Name, IdFacts ), + % fill it out + ( Mod:S, cuda_load_fact(S), fail ; true ). + +count_answers(G, N) :- + S = count(0), + ( + G, + arg(1, S, I0), + I is I0+1, + nb_setarg(1, S, I), + fail ; + S = count(N) ). cuda_rule((Head :- Body) , IdRules) :- body_to_list( Body, L, [], 1, N), diff --git a/packages/cuda/lista.cu b/packages/cuda/lista.cu index 4f4bb6131..f9f63bf9a 100644 --- a/packages/cuda/lista.cu +++ b/packages/cuda/lista.cu @@ -264,7 +264,7 @@ int2 columnsproject(int *first, int tam, int *rule, int ini, int fin, int sini, temp = first[x]; for(y = 0; y < ini; y++) { - if(temp == rule[y]) + if(temp == rule[y] && temp > 0) /*added condition to avoid constants*/ { if(notin(temp, pv, ret.y)) { @@ -279,7 +279,7 @@ int2 columnsproject(int *first, int tam, int *rule, int ini, int fin, int sini, continue; for(y = sfin + 1; y < fin; y++) { - if(temp == rule[y]) + if(temp == rule[y] && temp > 0) { if(notin(temp, pv, ret.y)) { @@ -298,7 +298,7 @@ int2 columnsproject(int *first, int tam, int *rule, int ini, int fin, int sini, temp = rule[x]; for(y = 0; y < ini; y++) { - if(temp == rule[y]) + if(temp == rule[y] && temp > 0) { if(notin(temp, pv, ret.y)) { @@ -313,7 +313,7 @@ int2 columnsproject(int *first, int tam, int *rule, int ini, int fin, int sini, continue; for(y = sfin + 1; y < fin; y++) { - if(temp == rule[y]) + if(temp == rule[y] && temp > 0) { if(notin(temp, pv, ret.y)) { @@ -1083,6 +1083,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, //cout << "resrows = " << res_rows << endl; } + if(x == num_refs) { if(rul_act->num_bpreds.x > 0) /*Built-in predicates*/ @@ -1094,7 +1095,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, cudaEventRecord(start3, 0); #endif - res_rows = bpreds(res, res_rows, rul_act->builtin, rul_act->num_bpreds, &res); + res_rows = bpreds(res, res_rows, rul_act->num_columns, rul_act->builtin, rul_act->num_bpreds, &res); #ifdef TIMER cudaEventRecord(stop3, 0); @@ -1115,7 +1116,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, cudaEventRecord(start2, 0); #endif - res_rows = unir(res, res_rows, rul_act->num_columns, &res); /*Duplicate Elimination*/ + res_rows = unir(res, res_rows, rul_act->num_columns); /*Duplicate Elimination*/ #ifdef TIMER cudaEventRecord(stop2, 0); @@ -1223,6 +1224,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, itr++; } + tmprule.name = qname; qposr = lower_bound(rul_str, fin, tmprule, comparer); if(qposr != fin && qposr->name == qname) @@ -1237,6 +1239,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, cols1 = qposf->num_columns; rows1 = cargar(qname, qposf->num_rows, cols1, 1, qposf->address_host_table, &dop1, 0); } + if(rows1 > 0) /*Query consideration*/ { consulta(query + 1, qsize, qname, &tmprule); @@ -1245,14 +1248,19 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, res = dop1; res_rows = rows1; } - else + else + { res_rows = selectproyect(dop1, rows1, cols1, tmprule.num_columns, tmprule.select[0], tmprule.numsel[0], tmprule.selfjoin[0], tmprule.numselfj[0], tmprule.project[0], &res); + if(qposr != fin && qposr->name == qname) + cudaFree(dop1); + } cols1 = tmprule.num_columns; tipo = res_rows * cols1 * sizeof(int); hres = (int *)malloc(tipo); cudaMemcpy(hres, res, tipo, cudaMemcpyDeviceToHost); - cudaFree(res); + if(res_rows > 0 && tmprule.numsel[0] != 0 && tmprule.numselfj[0] != 0) + cudaFree(res); } else res_rows = 0; @@ -1275,9 +1283,9 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, } //free(hres); - cout << "Elapsed = " << time << endl; - cout << "Size = " << res_rows << endl; - cout << "Iterations = " << itr << endl; + //cout << "Elapsed = " << time << endl; + //cout << "Size = " << res_rows << endl; + //cout << "Iterations = " << itr << endl; clear_memory(); *result = hres; diff --git a/packages/cuda/memory.cu b/packages/cuda/memory.cu index 413562396..933dddbd9 100644 --- a/packages/cuda/memory.cu +++ b/packages/cuda/memory.cu @@ -53,6 +53,7 @@ list::iterator buscarpornombre(int name, int itr, int *totalrows, int * { int x = 1, sum = 0; memnode temp; + temp.name = name; temp.iteration = itr; pair::iterator, list::iterator> rec = equal_range(GPUmem.begin(), GPUmem.end(), temp, compareiteration); @@ -104,14 +105,14 @@ int buscarpornombrecpu(int name, int itr, int *totalrows) return x; } -void limpiar() +void limpiar(const char s[]) { list::iterator ini; memnode temp; if(GPUmem.size() == 0) { - cerr << "Not enough GPU memory: have " << avmem << endl; + cerr << s << ": not enough GPU memory: have " << avmem << endl; exit(1); } @@ -182,12 +183,26 @@ void liberar(int *ptr, int size) void reservar(int **ptr, int size) { - // cout << "R " << avmem << " " << size + //size_t free, total; + //cudaMemGetInfo( &free, &total ); + // cerr << "R " << free << " " << size << endl; + if (size == 0) { + *ptr = NULL; + return; + } while(avmem < size) - limpiar(); + limpiar("not enough memory"); while(cudaMalloc(ptr, size) == cudaErrorMemoryAllocation) - limpiar(); + limpiar("error in memory allocation"); + if (! *ptr ) { + size_t free, total; + cudaMemGetInfo( &free, &total ); + cerr << "Could not allocate " << size << " bytes, only " << free << " avaliable from total of " << total << " !!!" << endl; + cerr << "Exiting CUDA...." << endl; + exit(1); + } + // cerr << *ptr << " " << size; avmem -= size; // cout << " " << avmem << endl; @@ -235,6 +250,11 @@ int numrows(int name, int itr) return sum; } + + extern "C" void * YAP_IntToAtom(int); + extern "C" char * YAP_AtomName(void *); + + int cargar(int name, int num_rows, int num_columns, int is_fact, int *address_host_table, int **ptr, int itr) { int numgpu, numcpu, totalrows = 0; @@ -242,6 +262,7 @@ int cargar(int name, int num_rows, int num_columns, int is_fact, int *address_ho int size, itrant; list::iterator i; memnode fact; + if(is_fact) { i = buscarhecho(GPUmem.begin(), GPUmem.end(), name); @@ -464,12 +485,16 @@ void clear_memory() { list::iterator ini; list::iterator fin; - ini = GPUmem.begin(); + ini = GPUmem.begin(); fin = GPUmem.end(); while(ini != fin) { - cudaFree(ini->dev_address); - ini++; + if (ini->isrule) { + cudaFree(ini->dev_address); + ini = GPUmem.erase(ini); + } else { + ini++; + } } ini = CPUmem.begin(); fin = CPUmem.end(); @@ -478,6 +503,5 @@ void clear_memory() free(ini->dev_address); ini++; } - GPUmem.clear(); CPUmem.clear(); } diff --git a/packages/cuda/memory.h b/packages/cuda/memory.h index 1a573c5d9..9cc08e331 100644 --- a/packages/cuda/memory.h +++ b/packages/cuda/memory.h @@ -11,7 +11,7 @@ using namespace std; void calcular_mem(int); void liberar(int*, int); -void limpiar(); +void limpiar(const char []); void limpiartodo(int*, int*); int cargar(int, int, int, int, int*, int**, int); int cargafinal(int, int, int**); diff --git a/packages/cuda/pred.h b/packages/cuda/pred.h index f8ebc5f80..e452de53c 100644 --- a/packages/cuda/pred.h +++ b/packages/cuda/pred.h @@ -11,6 +11,13 @@ typedef struct Nodo{ typedef gpunode predicate; +#define SBG_EQ (-1) +#define SBG_GT (-2) +#define SBG_LT (-3) +#define SBG_GE (-4) +#define SBG_LE (-5) +#define SBG_DF (-6) + int Cuda_Eval(predicate**, int, predicate**, int, predicate*, int**); #endif diff --git a/packages/cuda/treeb.cu b/packages/cuda/treeb.cu index 8a30c8fd4..60ca1ca8d 100755 --- a/packages/cuda/treeb.cu +++ b/packages/cuda/treeb.cu @@ -234,7 +234,7 @@ __global__ void gIndexMultiJoin(int *R, int *S, int g_locations[], int sLen, int { extern __shared__ int shared[]; int s_cur = blockIdx.x * blockDim.x + threadIdx.x; - int posr, poss, x, y, temp, ini; + int posr, poss, x, y, ini; if(threadIdx.x < wj) shared[threadIdx.x] = muljoin[threadIdx.x]; @@ -258,15 +258,17 @@ __global__ void gIndexMultiJoin(int *R, int *S, int g_locations[], int sLen, int poss = s_cur * of2; else poss = sloc[s_cur] * of2; - ini = r_cur - count; - for(x = 0; x < wj; x += 2) + ini = r_cur - count; + for(y = ini; y < r_cur; y++) { - posr = shared[x]; - temp = p2[poss + shared[x+1]]; - for(y = ini; y < r_cur; y++) + posr = mloc[y] * of1; + for(x = 0; x < wj; x += 2) { - if(p1[mloc[y] * of1 + posr] != temp) + if(p1[posr + shared[x]] != p2[poss + shared[x+1]]) + { count--; + break; + } } } if(count > 0) @@ -839,7 +841,7 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: } catch(std::bad_alloc &e) { - limpiar(); + limpiar("inclusive scan in join"); } } //thrust::inclusive_scan(res + 1, res + newLen, res + 1); @@ -1021,7 +1023,7 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: } catch(std::bad_alloc &e) { - limpiar(); + limpiar("inclusive scan in join"); } } @@ -1126,7 +1128,7 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: reservar(&d_Rout, resSize); if(numj > 2) { - cudaMemcpy(dcons + rule->num_columns, wherej + 2, muljoinsize, cudaMemcpyHostToDevice); + cudaMemcpy(dcons + projp.y, wherej + 2, muljoinsize, cudaMemcpyHostToDevice); multiJoinWithWrite<<>> (d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, projp.x, projp.y, posR, posS, muljoin); } else diff --git a/packages/cuda/union2.cu b/packages/cuda/union2.cu index 213a1619b..d6c1b62e0 100644 --- a/packages/cuda/union2.cu +++ b/packages/cuda/union2.cu @@ -77,7 +77,7 @@ struct o3 } }; -int unir(int *res, int rows, int tipo, int **ret) +int unir(int *res, int rows, int tipo) { thrust::device_ptr pt, re; thrust::device_ptr pt2, re2; @@ -102,7 +102,7 @@ int unir(int *res, int rows, int tipo, int **ret) } catch(std::bad_alloc &e) { - limpiar(); + limpiar("sort/unique in unir"); } } nrows = thrust::distance(pt, re); @@ -140,7 +140,7 @@ int unir(int *res, int rows, int tipo, int **ret) } catch(std::bad_alloc &e) { - limpiar(); + limpiar("sort/unique in unir"); } } nrows = thrust::distance(pt2, re2); @@ -179,7 +179,7 @@ int unir(int *res, int rows, int tipo, int **ret) } catch(std::bad_alloc &e) { - limpiar(); + limpiar("sort/unique in unir"); } } nrows = thrust::distance(pt3, re3); diff --git a/pl/threads.yap b/pl/threads.yap index c7345e7d6..774c7c20f 100644 --- a/pl/threads.yap +++ b/pl/threads.yap @@ -690,6 +690,7 @@ message_queue_property(Id, Prop) :- var(Term), !, '$do_error'(instantiation_error, Goal). '$check_message_queue_or_alias'(Term, Goal) :- + \+ integer(Term), \+ atom(Term), Term \= '$message_queue'(_), !, '$do_error'(domain_error(queue_or_alias, Term), Goal). @@ -697,7 +698,6 @@ message_queue_property(Id, Prop) :- \+ recorded('$queue', q(_,_,_,I,_), _), !, '$do_error'(existence_error(queue, '$message_queue'(I)), Goal). '$check_message_queue_or_alias'(Term, Goal) :- - atom(Term), \+ recorded('$queue', q(Term,_,_,_,_), _), !, '$do_error'(existence_error(queue, Term), Goal). '$check_message_queue_or_alias'(_, _). @@ -871,7 +871,7 @@ thread_property(Id, Prop) :- '$enumerate_threads'(Id) :- '$max_threads'(Max), Max1 is Max-1, - '$between'(0,Max1,Id), + between(0,Max1,Id), '$thread_stacks'(Id, _, _, _). '$thread_property'(Id, alias(Alias)) :-