diogo
/
yap-6.3
Archived
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge branch 'master' of ssh://git.dcc.fc.up.pt/yap-6.3

This commit is contained in:
Vitor Santos Costa 2013-10-16 10:40:12 +01:00
parent e19a910152 a01b319032
commit 83695cc8c8
14 changed files with 210 additions and 67 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

55
C/errors.c
View File

@ -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))

4
configure vendored
View File

@ -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"

2
configure.in
View File

@ -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"

3
library/dialect/swi/fli/swi.h
View File

@ -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;
}

4
packages/cuda/bpreds.cu
View File

@ -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;
}

76
packages/cuda/cuda.c
View File

@ -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);

18
packages/cuda/cuda.yap
View File

@ -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),

30
packages/cuda/lista.cu
View File

@ -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;

42
packages/cuda/memory.cu
View File

@ -53,6 +53,7 @@ list<memnode>::iterator buscarpornombre(int name, int itr, int *totalrows, int *
{
int x = 1, sum = 0;
memnode temp;
temp.name = name;
temp.iteration = itr;
pair<list<memnode>::iterator, list<memnode>::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<memnode>::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<memnode>::iterator i;
memnode fact;
if(is_fact)
{
i = buscarhecho(GPUmem.begin(), GPUmem.end(), name);
@ -464,12 +485,16 @@ void clear_memory()
{
list<memnode>::iterator ini;
list<memnode>::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();
}

2
packages/cuda/memory.h
View File

@ -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**);

7
packages/cuda/pred.h
View File

@ -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

22
packages/cuda/treeb.cu
View File

@ -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<rulenode>:
}
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<rulenode>:
}
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<rulenode>:
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<<<blockllen, numthreads, sizepro + muljoinsize>>> (d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, projp.x, projp.y, posR, posS, muljoin);
}
else

8
packages/cuda/union2.cu
View File

@ -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<int> pt, re;
thrust::device_ptr<s2> 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);

4
pl/threads.yap
View File

@ -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)) :-