// proto-G-typed-query.cu
// ---------------------------------------------------------------------------
// PROTOTYPE G: TYPED edges + VARIABLES = real Atomese Query (subgraph
// homomorphism), beyond untyped co-occurrence.
//
// Query (2 variables, 3 clauses, 2 predicates):
// find $X,$Y : (is-a $X animal) AND (eats $X $Y) AND (is-a $Y fish)
// = "animals that eat fish"
//
// Join plan (reuses Proto-A set-intersection as the inner op):
// candX = { X : is-a(X, animal) } (in-neighbours of 'animal' via is-a)
// candYB = { Y : is-a(Y, fish) } (sorted) (in-neighbours of 'fish' via is-a)
// answer = { (X,Y) : X in candX, Y in eats-out(X) ∩ candYB }
// GPU: one block per X, intersect eats-out(X) with candYB.
//
// Verifies: (a) known hand-built answer; (b) GPU==CPU on a random typed graph;
// (c) GOLD cross-check vs the real opencog pattern-matcher (a 2-var MeetLink).
// (This is ONE query shape — a real typed multi-var pattern — not yet a general
// engine; the general case is GSI/cuTS-style, future.)
//
// Build: nvcc -O2 -arch=sm_75 -o proto-G proto-G-typed-query.cu
// Run: ./proto-G (then: guile gold.scm | sort | diff - <(sort expected.txt))
// ---------------------------------------------------------------------------
#include <cstdio>
#include <cstdint>
#include <vector>
#include <string>
#include <map>
#include <set>
#include <algorithm>
#include <random>
#include <cuda_runtime.h>
#define CUDA_OK(c) do{cudaError_t e=(c); if(e!=cudaSuccess){\
printf("CUDA ERR %s:%d %s\n",__FILE__,__LINE__,cudaGetErrorString(e));return -1;}}while(0)
enum { P_ISA=0, P_EATS=1, NPRED=2 };
// Typed graph: per predicate, in-adjacency (by dst) and out-adjacency (by src).
struct Graph {
int V=0; std::vector<std::string> name; std::map<std::string,int> id;
std::map<int,std::set<int>> out_[NPRED]; // out_[p][src] = {dst}
std::map<int,std::set<int>> in_[NPRED]; // in_[p][dst] = {src}
int W(const std::string& s){ auto it=id.find(s); if(it!=id.end())return it->second;
int i=V++; id[s]=i; name.push_back(s); return i; }
void edge(int p,const std::string&a,const std::string&b){ int s=W(a),d=W(b);
out_[p][s].insert(d); in_[p][d].insert(s); }
std::vector<uint32_t> sorted(const std::set<int>& s){ std::vector<uint32_t> v(s.begin(),s.end());
std::sort(v.begin(),v.end()); return v; }
};
__device__ __forceinline__ int bs(const uint32_t* a,int n,uint32_t k){int lo=0,hi=n-1;
while(lo<=hi){int m=(lo+hi)>>1;uint32_t v=a[m];if(v==k)return m;if(v<k)lo=m+1;else hi=m-1;}return -1;}
// one block per candidate X; intersect eats-out(X) with candYB -> (X,Y) pairs
__global__ void typed_join(const uint32_t* outoff,const uint32_t* outdst,
const uint32_t* candX,int nX,const uint32_t* candYB,int nYB,
uint32_t* ox,uint32_t* oy,int st,int* cnt){
int xi=blockIdx.x; if(xi>=nX)return; uint32_t X=candX[xi];
const uint32_t* Yl=outdst+outoff[X]; int lenY=outoff[X+1]-outoff[X];
for(int t=threadIdx.x;t<lenY;t+=blockDim.x){ uint32_t Y=Yl[t];
if(bs(candYB,nYB,Y)>=0){ int i=atomicAdd(cnt,1); if(i<st){ox[i]=X;oy[i]=Y;} } }
}
// run the join on GPU for query (is-a $X A)(eats $X $Y)(is-a $Y B)
static int gpu_query(Graph& G,int A,int B,std::vector<std::pair<uint32_t,uint32_t>>& out){
// out-CSR for EATS over all V
std::vector<uint32_t> ooff(G.V+1,0),odst;
for(int s=0;s<G.V;s++){ auto it=G.out_[P_EATS].find(s);
if(it!=G.out_[P_EATS].end()){ auto v=G.sorted(it->second); for(auto d:v)odst.push_back(d); ooff[s+1]=ooff[s]+v.size(); }
else ooff[s+1]=ooff[s]; }
auto candX = G.in_[P_ISA].count(A)? G.sorted(G.in_[P_ISA][A]) : std::vector<uint32_t>{};
auto candYB= G.in_[P_ISA].count(B)? G.sorted(G.in_[P_ISA][B]) : std::vector<uint32_t>{};
if(candX.empty()||candYB.empty()){ out.clear(); return 0; }
int st=candX.size()* (G.V) + 16; // generous bound
uint32_t *d_oo,*d_od,*d_cx,*d_cy,*d_ox,*d_oy; int* d_c;
CUDA_OK(cudaMalloc(&d_oo,ooff.size()*4)); CUDA_OK(cudaMalloc(&d_od,std::max((size_t)1,odst.size())*4));
CUDA_OK(cudaMalloc(&d_cx,candX.size()*4)); CUDA_OK(cudaMalloc(&d_cy,candYB.size()*4));
CUDA_OK(cudaMalloc(&d_ox,st*4)); CUDA_OK(cudaMalloc(&d_oy,st*4)); CUDA_OK(cudaMalloc(&d_c,4));
CUDA_OK(cudaMemcpy(d_oo,ooff.data(),ooff.size()*4,cudaMemcpyHostToDevice));
if(odst.size())CUDA_OK(cudaMemcpy(d_od,odst.data(),odst.size()*4,cudaMemcpyHostToDevice));
CUDA_OK(cudaMemcpy(d_cx,candX.data(),candX.size()*4,cudaMemcpyHostToDevice));
CUDA_OK(cudaMemcpy(d_cy,candYB.data(),candYB.size()*4,cudaMemcpyHostToDevice));
CUDA_OK(cudaMemset(d_c,0,4));
typed_join<<<candX.size(),128>>>(d_oo,d_od,d_cx,(int)candX.size(),d_cy,(int)candYB.size(),d_ox,d_oy,st,d_c);
CUDA_OK(cudaDeviceSynchronize());
int c; CUDA_OK(cudaMemcpy(&c,d_c,4,cudaMemcpyDeviceToHost)); c=std::min(c,st);
std::vector<uint32_t> ox(c),oy(c);
if(c){CUDA_OK(cudaMemcpy(ox.data(),d_ox,c*4,cudaMemcpyDeviceToHost));CUDA_OK(cudaMemcpy(oy.data(),d_oy,c*4,cudaMemcpyDeviceToHost));}
out.clear(); for(int i=0;i<c;i++) out.push_back({ox[i],oy[i]});
cudaFree(d_oo);cudaFree(d_od);cudaFree(d_cx);cudaFree(d_cy);cudaFree(d_ox);cudaFree(d_oy);cudaFree(d_c);
return 0;
}
static std::vector<std::pair<uint32_t,uint32_t>> cpu_query(Graph& G,int A,int B){
std::vector<std::pair<uint32_t,uint32_t>> r;
if(!G.in_[P_ISA].count(A)||!G.in_[P_ISA].count(B))return r;
auto YB=G.in_[P_ISA][B];
for(int X:G.in_[P_ISA][A]){ auto it=G.out_[P_EATS].find(X); if(it==G.out_[P_EATS].end())continue;
for(int Y:it->second) if(YB.count(Y)) r.push_back({(uint32_t)X,(uint32_t)Y}); }
std::sort(r.begin(),r.end()); return r;
}
int main(){
printf("=== Prototype G: typed edges + variables (real Atomese Query) ===\n\n");
int fails=0;
// ---- (a) hand-built KG with KNOWN answer ----------------------------
Graph G;
auto isa=[&](const char*a,const char*b){G.edge(P_ISA,a,b);};
auto eat=[&](const char*a,const char*b){G.edge(P_EATS,a,b);};
isa("cat","animal");isa("dog","animal");isa("bear","animal");isa("otter","animal");isa("eagle","animal");
isa("salmon","fish");isa("trout","fish");isa("mouse","rodent");isa("grass","plant");
eat("cat","mouse");eat("cat","salmon");eat("bear","salmon");eat("bear","trout");
eat("otter","trout");eat("otter","salmon");eat("eagle","trout");eat("dog","grass");
int A=G.id["animal"],B=G.id["fish"];
std::vector<std::pair<uint32_t,uint32_t>> g; gpu_query(G,A,B,g); std::sort(g.begin(),g.end());
auto c=cpu_query(G,A,B);
printf("Query: (is-a $X animal) AND (eats $X $Y) AND (is-a $Y fish)\n GPU results (%zu):\n",g.size());
for(auto&p:g) printf(" %s eats %s\n",G.name[p.first].c_str(),G.name[p.second].c_str());
// known answer
std::set<std::pair<std::string,std::string>> exp={{"cat","salmon"},{"bear","salmon"},{"bear","trout"},
{"otter","trout"},{"otter","salmon"},{"eagle","trout"}};
std::set<std::pair<std::string,std::string>> got; for(auto&p:g)got.insert({G.name[p.first],G.name[p.second]});
if(got==exp&&g==c)printf(" PASS (matches known answer AND GPU==CPU)\n");
else{fails++;printf(" FAIL known=%d gpu==cpu=%d\n",got==exp,g==c);}
// ---- emit gold.scm + expected.txt (real opencog cross-check) --------
FILE* scm=fopen("gold-G.scm","w"); FILE* ex=fopen("expected-G.txt","w");
fprintf(scm,"(use-modules (opencog) (opencog exec))\n");
const char* pn[NPRED]={"is-a","eats"};
for(int p=0;p<NPRED;p++) for(auto&kv:G.out_[p]) for(int d:kv.second)
fprintf(scm,"(Evaluation (Predicate \"%s\") (List (Concept \"%s\") (Concept \"%s\")))\n",
pn[p],G.name[kv.first].c_str(),G.name[d].c_str());
fprintf(scm,
"(define r (cog-execute! (Meet\n"
" (VariableList (TypedVariable (Variable \"$X\")(Type 'ConceptNode))\n"
" (TypedVariable (Variable \"$Y\")(Type 'ConceptNode)))\n"
" (And (Evaluation (Predicate \"is-a\") (List (Variable \"$X\")(Concept \"animal\")))\n"
" (Evaluation (Predicate \"eats\") (List (Variable \"$X\")(Variable \"$Y\")))\n"
" (Evaluation (Predicate \"is-a\") (List (Variable \"$Y\")(Concept \"fish\")))))))\n"
"(for-each (lambda (g) (format #t \"~a ~a\\n\" (cog-name (car (cog-outgoing-set g)))\n"
" (cog-name (cadr (cog-outgoing-set g)))))\n"
" (sort (cog-outgoing-set r)\n"
" (lambda (p q) (string<? (cog-name (car (cog-outgoing-set p)))\n"
" (cog-name (car (cog-outgoing-set q)))))))\n");
for(auto&p:g) fprintf(ex,"%s %s\n",G.name[p.first].c_str(),G.name[p.second].c_str());
fclose(scm);fclose(ex);
printf(" wrote gold-G.scm + expected-G.txt for opencog cross-check\n");
// ---- (b) random typed graph: GPU == CPU ----------------------------
Graph R; std::mt19937 rng(7);
const int ITEMS=4000, TYPES=40;
for(int t=0;t<TYPES;t++) R.W("type"+std::to_string(t));
for(int i=0;i<ITEMS;i++){ std::string it="it"+std::to_string(i);
int nt=1+rng()%2; for(int k=0;k<nt;k++) R.edge(P_ISA,it,"type"+std::to_string(rng()%TYPES)); }
for(int i=0;i<ITEMS;i++){ std::string it="it"+std::to_string(i); int ne=rng()%6;
for(int k=0;k<ne;k++) R.edge(P_EATS,it,"it"+std::to_string(rng()%ITEMS)); }
int rfail=0,RT=200;
for(int q=0;q<RT;q++){ int a=R.id["type"+std::to_string(rng()%TYPES)],b=R.id["type"+std::to_string(rng()%TYPES)];
std::vector<std::pair<uint32_t,uint32_t>> rg; gpu_query(R,a,b,rg); std::sort(rg.begin(),rg.end());
auto rc=cpu_query(R,a,b); if(rg!=rc)rfail++; }
printf("\n(b) random typed graph: %d items, %d types, %d queries -> GPU==CPU on %d/%d\n",
ITEMS,TYPES,RT,RT-rfail,RT); if(rfail){fails++;printf(" FAIL %d mismatches\n",rfail);} else printf(" PASS\n");
printf("\n%s\n", fails? "SOME CHECKS FAILED" :
"ALL CHECKS PASSED - typed multi-variable Query works on GPU.");
return fails?1:0;
}