diff --git a/packages/cuda/Makefile.in b/packages/cuda/Makefile.in index 251194261..8690e6840 100644 --- a/packages/cuda/Makefile.in +++ b/packages/cuda/Makefile.in @@ -53,7 +53,7 @@ all: $(SOBJS) cuda.o: $(srcdir)/cuda.c $(srcdir)/pred.h $(NVCC) -c $(NVCCFLAGS) $(srcdir)/cuda.c -o cuda.o -lista.o: $(srcdir)/lista.cu $(srcdir)/pred.h +lista.o: $(srcdir)/lista.cu $(srcdir)/pred.h $(srcdir)/selectproyect.cu $(srcdir)/treeb.cu $(srcdir)/union2.cu $(srcdir)/bpreds.cu $(NVCC) -c $(NVCCFLAGS) $(srcdir)/lista.cu -o lista.o memory.o: $(srcdir)/memory.cu $(srcdir)/pred.h diff --git a/packages/cuda/bpreds.cu b/packages/cuda/bpreds.cu index ea7814980..ea94b5b2a 100644 --- a/packages/cuda/bpreds.cu +++ b/packages/cuda/bpreds.cu @@ -1,10 +1,3 @@ -#define SBG_EQ (-1) -#define SBG_GT (-2) -#define SBG_LT (-3) -#define SBG_GE (-4) -#define SBG_LE (-5) -#define SBG_DF (-6) - __global__ void predicates(int *dop1, int rows, int cols, int *cons, int numc, int *res) { extern __shared__ int shared[]; @@ -30,16 +23,21 @@ __global__ void predicates(int *dop1, int rows, int cols, int *cons, int numc, i op2 = dop1[rowact + op2]; switch(shared[x]) { - case SBG_EQ: if(op1 != op2) + case SBG_EQ: if(op1 != op2) return; + break; case SBG_GT: if(op1 <= op2) return; + break; case SBG_LT: if(op1 >= op2) return; + break; case SBG_GE: if(op1 < op2) return; + break; case SBG_LE: if(op1 > op2) return; + break; case SBG_DF: if(op1 == op2) return; } diff --git a/packages/cuda/cuda.c b/packages/cuda/cuda.c index 6929b0bd3..7684b7c10 100644 --- a/packages/cuda/cuda.c +++ b/packages/cuda/cuda.c @@ -8,6 +8,13 @@ #include #include "pred.h" +YAP_Atom AtomEq, + AtomGt, + AtomLt, + AtomGe, + AtomLe, + AtomDf; + predicate *facts[100]; /*Temporary solution to maintain facts and rules*/ predicate *rules[100]; int32_t cf = 0, cr = 0; @@ -160,8 +167,22 @@ load_rule( void ) { YAP_Term th = YAP_HeadOfTerm(t3); YAP_Functor f = YAP_FunctorOfTerm( th ); int32_t n = YAP_ArityOfFunctor( f ); + YAP_Atom at = YAP_NameOfFunctor( f ); - *ptr++ = YAP_AtomToInt( YAP_NameOfFunctor( f ) ); + if (at == AtomEq) + *ptr++ = SBG_EQ; + else if (at == AtomGt) + *ptr++ = SBG_GT; + else if (at == AtomLt) + *ptr++ = SBG_LT; + else if (at == AtomGe) + *ptr++ = SBG_GE; + else if (at == AtomLe) + *ptr++ = SBG_LE; + else if (at == AtomDf) + *ptr++ = SBG_DF; + else + *ptr++ = YAP_AtomToInt( at ); for (j = 0; j < n; j++) { YAP_Term ta = YAP_ArgOfTerm(j+1, th); @@ -257,6 +278,12 @@ init_cuda(void) if (first_time) Cuda_Initialize(); first_time = FALSE; + AtomEq = YAP_LookupAtom("="); + AtomGt = YAP_LookupAtom(">"); + AtomLt = YAP_LookupAtom("<"); + AtomGe = YAP_LookupAtom(">="); + AtomLe = YAP_LookupAtom("=<"); + AtomDf = YAP_LookupAtom("\\="); YAP_UserCPredicate("load_facts", load_facts, 4); YAP_UserCPredicate("load_rule", load_rule, 4); YAP_UserCPredicate("cuda_erase", cuda_erase, 1); diff --git a/packages/cuda/test.yap b/packages/cuda/test.yap index b5c17cd47..82c49f768 100644 --- a/packages/cuda/test.yap +++ b/packages/cuda/test.yap @@ -1,12 +1,15 @@ :- use_module(library(cuda)). +:- use_module(library(lists)). :- initialization(main). main :- - cuda_extensional(db/2, _X), - cuda_rule((a(X, Y) :- db(Y, Z), db(X, Z), db(1, Z) ), Q), - cuda_eval(Q, L), writeln(here), + Rule = ( db(Y, Z), db(X, Z), db(1, Z), X = Y ), + setof(a(X,Y), Z^Rule, L0), reverse(L0, RL0), writeln(RL0), + cuda_rule((a(X, Y) :- Rule ), Q), + cuda_eval(Q, L), + cuda_erase( Q ), writeln(L). db(1,a). @@ -15,3 +18,19 @@ db(5,b). db(4,q). db(6,w). db(10,s). +/* +db(11,a). +db(12,a). +db(15,b). +db(14,q). +db(16,w). +db(110,s). +db(21,a). +db(22,a). +db(25,b). +db(24,q). +db(26,w). +db(210,s). +*/ + +:- cuda_extensional(db/2, _X). \ No newline at end of file