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

more CUDA fixes

This commit is contained in:
Vítor Santos Costa 2013-10-07 13:48:19 +01:00
parent 1593ee4918
commit fbf6648433
4 changed files with 57 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
packages/cuda/Makefile.in
View File

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

14
packages/cuda/bpreds.cu
View File

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

29
packages/cuda/cuda.c
View File

@ -8,6 +8,13 @@
#include <string.h>
#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);

25
packages/cuda/test.yap
View File

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