fecc49c · atomspace-gpu

Changed files (8) +859 -2

opencog/persist/gpu/GpuAtomTable.cu +151 -0

…	…	@@ -17,6 +17,7 @@
17	17	#include <cuda_runtime.h>
18	18	#include <cstdio>
19	19	#include <cstring>
	20	#include <vector>
20	21
21	22	/* ---------------------------------------------------------------- */
22	23	#define CUDA_CHECK(call, ctx) do { \
…	…	@@ -50,6 +51,10 @@ int cuda_table_alloc(GpuAtomTable* t)
50	51	"alloc name_pool");
51	52	CUDA_CHECK(cudaMalloc(&t->out_pool, GPU_OUT_POOL_SLOTS * sizeof(uint32_t)),
52	53	"alloc out_pool");
	54	CUDA_CHECK(cudaMalloc(&t->in_off, (GPU_ATOM_CAPACITY + 1) * sizeof(uint32_t)),
	55	"alloc in_off");
	56	CUDA_CHECK(cudaMalloc(&t->in_pool, GPU_IN_POOL_SLOTS * sizeof(uint32_t)),
	57	"alloc in_pool");
53	58
54	59	CUDA_CHECK(cudaMemset(t->atoms, 0, GPU_ATOM_CAPACITY * sizeof(GpuAtom)),
55	60	"zero atoms");
…	…	@@ -57,10 +62,14 @@ int cuda_table_alloc(GpuAtomTable* t)
57	62	"zero name_pool");
58	63	CUDA_CHECK(cudaMemset(t->out_pool, 0, GPU_OUT_POOL_SLOTS * sizeof(uint32_t)),
59	64	"zero out_pool");
	65	CUDA_CHECK(cudaMemset(t->in_off, 0, (GPU_ATOM_CAPACITY + 1) * sizeof(uint32_t)),
	66	"zero in_off");
60	67
61	68	t->atom_count = 0;
62	69	t->name_pool_used = 0;
63	70	t->out_pool_used = 0;
	71	t->in_pool_used = 0;
	72	t->in_built = 0;
64	73
65	74	return 0;
66	75	}
…	…	@@ -74,6 +83,8 @@ void cuda_table_free(GpuAtomTable* t)
74	83	if (t->atoms) cudaFree(t->atoms);
75	84	if (t->name_pool) cudaFree(t->name_pool);
76	85	if (t->out_pool) cudaFree(t->out_pool);
	86	if (t->in_off) cudaFree(t->in_off);
	87	if (t->in_pool) cudaFree(t->in_pool);
77	88
78	89	memset(t, 0, sizeof(GpuAtomTable));
79	90	}
…	…	@@ -90,10 +101,14 @@ void cuda_table_clear(GpuAtomTable* t)
90	101	"clear name_pool");
91	102	CUDA_CHECK_VOID(cudaMemset(t->out_pool, 0, GPU_OUT_POOL_SLOTS * sizeof(uint32_t)),
92	103	"clear out_pool");
	104	CUDA_CHECK_VOID(cudaMemset(t->in_off, 0, (GPU_ATOM_CAPACITY + 1) * sizeof(uint32_t)),
	105	"clear in_off");
93	106
94	107	t->atom_count = 0;
95	108	t->name_pool_used = 0;
96	109	t->out_pool_used = 0;
	110	t->in_pool_used = 0;
	111	t->in_built = 0;
97	112	}
98	113
99	114	/* ================================================================
…	…	@@ -230,3 +245,139 @@ void cuda_table_barrier(GpuAtomTable* t)
230	245	(void)t;
231	246	cudaDeviceSynchronize();
232	247	}
	248
	249	/* ================================================================
	250	* Incoming-set index (Phase 2a).
	251	* Build a CSR of link-slots referencing each atom, from the stored
	252	* outgoing sets. count -> exclusive-scan -> scatter.
	253	* ================================================================ */
	254
	255	__global__ void k_count_incoming(const GpuAtom* atoms, const uint32_t* out_pool,
	256	uint32_t n, uint32_t* counts)
	257	{
	258	uint32_t s = blockIdx.x * blockDim.x + threadIdx.x;
	259	if (s >= n) return;
	260	GpuAtom a = atoms[s];
	261	if (a.is_node) return;
	262	for (uint16_t i = 0; i < a.data_len; i++)
	263	atomicAdd(&counts[out_pool[a.data_offset + i]], 1u);
	264	}
	265
	266	__global__ void k_scatter_incoming(const GpuAtom* atoms, const uint32_t* out_pool,
	267	uint32_t n, uint32_t* cursor, uint32_t* in_pool)
	268	{
	269	uint32_t s = blockIdx.x * blockDim.x + threadIdx.x;
	270	if (s >= n) return;
	271	GpuAtom a = atoms[s];
	272	if (a.is_node) return;
	273	for (uint16_t i = 0; i < a.data_len; i++) {
	274	uint32_t o = out_pool[a.data_offset + i];
	275	uint32_t pos = atomicAdd(&cursor[o], 1u);
	276	in_pool[pos] = s;
	277	}
	278	}
	279
	280	extern "C"
	281	int cuda_table_build_incoming(GpuAtomTable* t)
	282	{
	283	uint32_t n = t->atom_count;
	284	t->in_built = 0;
	285	t->in_pool_used = 0;
	286	if (n == 0) { t->in_built = 1; return 0; }
	287
	288	uint32_t* d_counts;
	289	CUDA_CHECK(cudaMalloc(&d_counts, n * sizeof(uint32_t)), "alloc counts");
	290	CUDA_CHECK(cudaMemset(d_counts, 0, n * sizeof(uint32_t)), "zero counts");
	291
	292	int TB = 256;
	293	int GB = (int)((n + TB - 1) / TB);
	294	k_count_incoming<<<GB, TB>>>(t->atoms, t->out_pool, n, d_counts);
	295	CUDA_CHECK(cudaDeviceSynchronize(), "count incoming");
	296
	297	std::vector<uint32_t> counts(n);
	298	CUDA_CHECK(cudaMemcpy(counts.data(), d_counts, n * sizeof(uint32_t),
	299	cudaMemcpyDeviceToHost), "copy counts");
	300	std::vector<uint32_t> off(n + 1);
	301	off[0] = 0;
	302	for (uint32_t i = 0; i < n; i++) off[i + 1] = off[i] + counts[i];
	303	uint32_t total = off[n];
	304	if (total > GPU_IN_POOL_SLOTS) { cudaFree(d_counts); return -1; }
	305	CUDA_CHECK(cudaMemcpy(t->in_off, off.data(), (n + 1) * sizeof(uint32_t),
	306	cudaMemcpyHostToDevice), "upload in_off");
	307
	308	uint32_t* d_cursor;
	309	CUDA_CHECK(cudaMalloc(&d_cursor, n * sizeof(uint32_t)), "alloc cursor");
	310	CUDA_CHECK(cudaMemcpy(d_cursor, off.data(), n * sizeof(uint32_t),
	311	cudaMemcpyHostToDevice), "init cursor");
	312	k_scatter_incoming<<<GB, TB>>>(t->atoms, t->out_pool, n, d_cursor, t->in_pool);
	313	CUDA_CHECK(cudaDeviceSynchronize(), "scatter incoming");
	314
	315	cudaFree(d_counts);
	316	cudaFree(d_cursor);
	317	t->in_pool_used = total;
	318	t->in_built = 1;
	319	return 0;
	320	}
	321
	322	extern "C"
	323	int cuda_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
	324	uint32_t* buf, uint32_t max, uint32_t* count)
	325	{
	326	*count = 0;
	327	if (slot >= t->atom_count) return -1;
	328
	329	uint32_t bounds[2];
	330	CUDA_CHECK(cudaMemcpy(bounds, t->in_off + slot, 2 * sizeof(uint32_t),
	331	cudaMemcpyDeviceToHost), "fetch in_off");
	332	uint32_t c = bounds[1] - bounds[0];
	333	*count = c;
	334	uint32_t k = (c < max) ? c : max;
	335	if (k > 0 && buf)
	336	CUDA_CHECK(cudaMemcpy(buf, t->in_pool + bounds[0], k * sizeof(uint32_t),
	337	cudaMemcpyDeviceToHost), "fetch incoming");
	338	return 0;
	339	}
	340
	341	/* GPU-resident join: intersect incoming(a) and incoming(b) on the GPU. */
	342	__global__ void k_intersect_incoming(const uint32_t* in_off, const uint32_t* in_pool,
	343	uint32_t a, uint32_t b,
	344	uint32_t* out, int maxn, int* cnt)
	345	{
	346	uint32_t as = in_off[a], ae = in_off[a + 1];
	347	uint32_t bs = in_off[b], be = in_off[b + 1];
	348	for (uint32_t i = as + threadIdx.x; i < ae; i += blockDim.x) {
	349	uint32_t x = in_pool[i];
	350	bool found = false;
	351	for (uint32_t j = bs; j < be; j++) if (in_pool[j] == x) { found = true; break; }
	352	if (found) { int k = atomicAdd(cnt, 1); if (k < maxn) out[k] = x; }
	353	}
	354	}
	355
	356	extern "C"
	357	int cuda_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
	358	uint32_t* outbuf, uint32_t max, uint32_t* count)
	359	{
	360	*count = 0;
	361	if (a >= t->atom_count \|\| b >= t->atom_count) return -1;
	362
	363	int* d_cnt;
	364	uint32_t* d_out;
	365	uint32_t alloc = (max > 0) ? max : 1;
	366	CUDA_CHECK(cudaMalloc(&d_cnt, sizeof(int)), "alloc isect cnt");
	367	CUDA_CHECK(cudaMemset(d_cnt, 0, sizeof(int)), "zero isect cnt");
	368	CUDA_CHECK(cudaMalloc(&d_out, alloc * sizeof(uint32_t)), "alloc isect out");
	369
	370	k_intersect_incoming<<<1, 128>>>(t->in_off, t->in_pool, a, b, d_out, (int)max, d_cnt);
	371	CUDA_CHECK(cudaDeviceSynchronize(), "intersect incoming");
	372
	373	int c;
	374	CUDA_CHECK(cudaMemcpy(&c, d_cnt, sizeof(int), cudaMemcpyDeviceToHost), "copy isect cnt");
	375	*count = (uint32_t)c;
	376	uint32_t k = ((uint32_t)c < max) ? (uint32_t)c : max;
	377	if (k > 0 && outbuf)
	378	CUDA_CHECK(cudaMemcpy(outbuf, d_out, k * sizeof(uint32_t),
	379	cudaMemcpyDeviceToHost), "copy isect out");
	380	cudaFree(d_cnt);
	381	cudaFree(d_out);
	382	return 0;
	383	}

opencog/persist/gpu/GpuAtomTable.h +39 -0

…	…	@@ -49,6 +49,13 @@
49	49	#define GPU_OUT_POOL_SLOTS 4194304 /* 4M outgoing-set entries */
50	50	#endif
51	51
	52	/* Incoming-set index (CSR): total incoming entries == total outgoing
	53	* entries (each outgoing reference creates exactly one incoming entry),
	54	* so the incoming pool is sized to match the outgoing pool. */
	55	#ifndef GPU_IN_POOL_SLOTS
	56	#define GPU_IN_POOL_SLOTS 4194304 /* 4M incoming entries */
	57	#endif
	58
52	59	/* -----------------------------------------------------------
53	60	* GpuBackendType -- which GPU backend is active for a table.
54	61	* ----------------------------------------------------------- */
…	…	@@ -90,10 +97,17 @@ typedef struct GpuAtomTable
90	97	char* name_pool; /* [NAME_POOL_BYTES] packed node names */
91	98	uint32_t* out_pool; /* [OUT_POOL_SLOTS] packed outgoing sets */
92	99
	100	/* Incoming-set CSR index (built on demand from links). For atom slot s,
	101	* its incoming links are in_pool[in_off[s] .. in_off[s+1]). */
	102	uint32_t* in_off; /* [CAPACITY+1] device CSR offsets */
	103	uint32_t* in_pool; /* [IN_POOL_SLOTS] device incoming link-slots */
	104
93	105	/* Host-side bookkeeping (NOT on device) */
94	106	uint32_t atom_count; /* next free slot index */
95	107	uint32_t name_pool_used; /* bytes used in name_pool */
96	108	uint32_t out_pool_used; /* slots used in out_pool */
	109	uint32_t in_pool_used; /* entries used in in_pool (after build) */
	110	uint8_t in_built; /* 1 once incoming index is current */
97	111	} GpuAtomTable;
98	112
99	113	/* -----------------------------------------------------------
…	…	@@ -123,6 +137,21 @@ int gpu_fetch_node(const GpuAtomTable* t, uint32_t slot,
123	137	int gpu_fetch_link(const GpuAtomTable* t, uint32_t slot,
124	138	uint16_t* type, uint32_t* out_buf, uint16_t* arity);
125	139
	140	/* Build the incoming-set CSR index from the stored links. Call after
	141	* storing atoms and before querying incoming sets. Idempotent. */
	142	int gpu_table_build_incoming(GpuAtomTable* t);
	143
	144	/* Fetch the incoming set (link-slots referencing `slot`) into buf (up to
	145	* max entries). Writes the true count to count (may exceed max). /
	146	int gpu_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
	147	uint32_t* buf, uint32_t max, uint32_t* count);
	148
	149	/* GPU-resident join: intersect the incoming sets of atoms `a` and `b`
	150	* (links referencing BOTH) entirely on the GPU. Requires the incoming
	151	* index to be built. */
	152	int gpu_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
	153	uint32_t* out, uint32_t max, uint32_t* count);
	154
126	155	/* Synchronization */
127	156	void gpu_table_barrier(GpuAtomTable* t);
128	157
…	…	@@ -143,6 +172,11 @@ int cuda_fetch_node(const GpuAtomTable* t, uint32_t slot,
143	172	uint16_t* type, char* name_buf, uint16_t* name_len);
144	173	int cuda_fetch_link(const GpuAtomTable* t, uint32_t slot,
145	174	uint16_t* type, uint32_t* out_buf, uint16_t* arity);
	175	int cuda_table_build_incoming(GpuAtomTable* t);
	176	int cuda_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
	177	uint32_t* buf, uint32_t max, uint32_t* count);
	178	int cuda_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
	179	uint32_t* out, uint32_t max, uint32_t* count);
146	180	void cuda_table_barrier(GpuAtomTable* t);
147	181	#endif
148	182
…	…	@@ -159,6 +193,11 @@ int ocl_fetch_node(const GpuAtomTable* t, uint32_t slot,
159	193	uint16_t* type, char* name_buf, uint16_t* name_len);
160	194	int ocl_fetch_link(const GpuAtomTable* t, uint32_t slot,
161	195	uint16_t* type, uint32_t* out_buf, uint16_t* arity);
	196	int ocl_table_build_incoming(GpuAtomTable* t);
	197	int ocl_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
	198	uint32_t* buf, uint32_t max, uint32_t* count);
	199	int ocl_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
	200	uint32_t* out, uint32_t max, uint32_t* count);
162	201	void ocl_table_barrier(GpuAtomTable* t);
163	202	#endif
164	203

opencog/persist/gpu/GpuAtomTable_dispatch.cc +56 -0

…	…	@@ -186,3 +186,59 @@ void gpu_table_barrier(GpuAtomTable* t)
186	186	default: break;
187	187	}
188	188	}
	189
	190	/* ================================================================
	191	* gpu_table_build_incoming
	192	* ================================================================ */
	193	extern "C"
	194	int gpu_table_build_incoming(GpuAtomTable* t)
	195	{
	196	switch (t->backend)
	197	{
	198	#ifdef HAVE_CUDA
	199	case GPU_BACKEND_CUDA: return cuda_table_build_incoming(t);
	200	#endif
	201	#ifdef HAVE_OPENCL
	202	case GPU_BACKEND_OPENCL: return ocl_table_build_incoming(t);
	203	#endif
	204	default: return -1;
	205	}
	206	}
	207
	208	/* ================================================================
	209	* gpu_fetch_incoming
	210	* ================================================================ */
	211	extern "C"
	212	int gpu_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
	213	uint32_t* buf, uint32_t max, uint32_t* count)
	214	{
	215	switch (t->backend)
	216	{
	217	#ifdef HAVE_CUDA
	218	case GPU_BACKEND_CUDA: return cuda_fetch_incoming(t, slot, buf, max, count);
	219	#endif
	220	#ifdef HAVE_OPENCL
	221	case GPU_BACKEND_OPENCL: return ocl_fetch_incoming(t, slot, buf, max, count);
	222	#endif
	223	default: return -1;
	224	}
	225	}
	226
	227	/* ================================================================
	228	* gpu_intersect_incoming
	229	* ================================================================ */
	230	extern "C"
	231	int gpu_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
	232	uint32_t* out, uint32_t max, uint32_t* count)
	233	{
	234	switch (t->backend)
	235	{
	236	#ifdef HAVE_CUDA
	237	case GPU_BACKEND_CUDA: return cuda_intersect_incoming(t, a, b, out, max, count);
	238	#endif
	239	#ifdef HAVE_OPENCL
	240	case GPU_BACKEND_OPENCL: return ocl_intersect_incoming(t, a, b, out, max, count);
	241	#endif
	242	default: return -1;
	243	}
	244	}

opencog/persist/gpu/GpuAtomTable_ocl.cc +31 -0

…	…	@@ -340,4 +340,35 @@ void ocl_table_barrier(GpuAtomTable* t)
340	340	}
341	341	}
342	342
	343	/* ================================================================
	344	* Incoming-set index (Phase 2a) — not yet implemented for OpenCL.
	345	* The CUDA backend is preferred; these stubs keep the dispatch
	346	* layer linkable. Returns -1 (caller treats as "no incoming").
	347	* ================================================================ */
	348	extern "C"
	349	int ocl_table_build_incoming(GpuAtomTable* t)
	350	{
	351	(void)t;
	352	fprintf(stderr, "OpenCL: build_incoming not implemented (use CUDA)\n");
	353	return -1;
	354	}
	355
	356	extern "C"
	357	int ocl_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
	358	uint32_t* buf, uint32_t max, uint32_t* count)
	359	{
	360	(void)t; (void)slot; (void)buf; (void)max;
	361	if (count) *count = 0;
	362	return -1;
	363	}
	364
	365	extern "C"
	366	int ocl_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
	367	uint32_t* out, uint32_t max, uint32_t* count)
	368	{
	369	(void)t; (void)a; (void)b; (void)out; (void)max;
	370	if (count) *count = 0;
	371	return -1;
	372	}
	373
343	374	#endif /* HAVE_OPENCL */

opencog/persist/gpu/GpuStorageNode.cc +70 -2

…	…	@@ -152,6 +152,9 @@ void GpuStorageNode::do_store_atom(const Handle& h)
152	152	if (rc != 0)
153	153	throw IOException(TRACE_INFO,
154	154	"GPU: failed to store atom at slot %u", slot);
	155
	156	// New atom invalidates the incoming-set index.
	157	_gpu_table.in_built = 0;
155	158	}
156	159
157	160	void GpuStorageNode::storeAtom(const Handle& h, bool synchronous)
…	…	@@ -236,8 +239,73 @@ void GpuStorageNode::getAtom(const Handle&)
236	239	// Step 3: no Values on GPU. Nothing to load.
237	240	}
238	241
239		void GpuStorageNode::fetchIncomingSet(AtomSpace*, const Handle&) {}
240		void GpuStorageNode::fetchIncomingByType(AtomSpace*, const Handle&, Type) {}
	242	// Build the incoming-set CSR index on the GPU if it is stale.
	243	void GpuStorageNode::ensure_incoming()
	244	{
	245	if (!_gpu_table.in_built)
	246	gpu_table_build_incoming(&_gpu_table);
	247	}
	248
	249	// Reconstruct one GPU link-slot as a Link in `as`. Returns the Handle
	250	// (or undefined if any outgoing atom is unknown). type_filter == NOTYPE
	251	// means "any type".
	252	Handle GpuStorageNode::fetch_link_to(AtomSpace* as, uint32_t link_slot,
	253	Type type_filter)
	254	{
	255	uint16_t ltype = 0;
	256	uint32_t out_buf[256];
	257	uint16_t arity = sizeof(out_buf) / sizeof(out_buf[0]);
	258	if (gpu_fetch_link(&_gpu_table, link_slot, &ltype, out_buf, &arity) != 0)
	259	return Handle::UNDEFINED;
	260	if (NOTYPE != type_filter && (Type)ltype != type_filter)
	261	return Handle::UNDEFINED;
	262
	263	HandleSeq oset;
	264	oset.reserve(arity);
	265	for (uint16_t i = 0; i < arity; i++)
	266	{
	267	Handle oh = _tlb.getAtom((UUID)out_buf[i]);
	268	if (nullptr == oh) return Handle::UNDEFINED;
	269	oset.push_back(oh);
	270	}
	271	return as->add_link((Type)ltype, std::move(oset));
	272	}
	273
	274	// Fetch the incoming set of `h` (optionally filtered by type) from the GPU
	275	// and install the links into `as`.
	276	void GpuStorageNode::fetch_incoming_impl(AtomSpace* as, const Handle& h,
	277	Type type_filter)
	278	{
	279	std::lock_guard<std::mutex> lk(_mtx);
	280	if (!_is_open) return;
	281
	282	auto it = _handle_to_slot.find(h);
	283	if (it == _handle_to_slot.end()) return; // atom not on GPU
	284
	285	ensure_incoming();
	286
	287	uint32_t slot = it->second;
	288	uint32_t count = 0;
	289	gpu_fetch_incoming(&_gpu_table, slot, nullptr, 0, &count);
	290	if (0 == count) return;
	291
	292	std::vector<uint32_t> links(count);
	293	gpu_fetch_incoming(&_gpu_table, slot, links.data(), count, &count);
	294
	295	for (uint32_t ls : links)
	296	fetch_link_to(as, ls, type_filter);
	297	}
	298
	299	void GpuStorageNode::fetchIncomingSet(AtomSpace* as, const Handle& h)
	300	{
	301	fetch_incoming_impl(as, h, NOTYPE);
	302	}
	303
	304	void GpuStorageNode::fetchIncomingByType(AtomSpace* as, const Handle& h, Type t)
	305	{
	306	fetch_incoming_impl(as, h, t);
	307	}
	308
241	309	void GpuStorageNode::removeAtom(AtomSpace*, const Handle&, bool) {}
242	310	void GpuStorageNode::loadType(AtomSpace*, Type) {}
243	311

opencog/persist/gpu/GpuStorageNode.h +3 -0

…	…	@@ -47,6 +47,9 @@ private:
47	47	bool has_slot(const Handle&) const;
48	48	uint32_t assign_slot(const Handle&);
49	49	void do_store_atom(const Handle&);
	50	void ensure_incoming(); // build incoming-set index if stale
	51	Handle fetch_link_to(AtomSpace*, uint32_t link_slot, Type type_filter);
	52	void fetch_incoming_impl(AtomSpace*, const Handle&, Type type_filter);
50	53
51	54	public:
52	55	GpuStorageNode(Type t, const std::string&& uri);

tests/persist/gpu/GpuBasicSaveUTest.cxxtest +430 -0

…	…	@@ -24,7 +24,10 @@
24	24	#include <opencog/persist/gpu/GpuStorageNode.h>
25	25	#include <opencog/persist/gpu/GpuAtomTable.h>
26	26	#include <opencog/persist/gpu-types/atom_types.h>
	27	#include <opencog/atoms/value/LinkValue.h>
27	28
	29	#include <set>
	30	#include <string>
28	31	#include <opencog/util/Logger.h>
29	32
30	33	using namespace opencog;
…	…	@@ -49,8 +52,77 @@ public:
49	52	void test_store_nodes(void);
50	53	void test_store_link(void);
51	54	void test_roundtrip(void);
	55	void test_fetch_incoming(void);
	56	void test_query(void);
	57	void test_v1_multihop(void);
	58	void test_v4_neurosymbolic(void);
	59	void test_agent_loop(void);
52	60	};
53	61
	62	// Execute a 3-variable MeetLink; return (X,Z) pairs = (outgoing[0], outgoing[2]).
	63	static std::vector<std::pair<Handle,Handle>> run_xz(AtomSpace* as, Handle vd, Handle body)
	64	{
	65	Handle meet = as->add_link(MEET_LINK, vd, body);
	66	ValuePtr sv = meet->execute(as);
	67	std::vector<std::pair<Handle,Handle>> out;
	68	LinkValuePtr lv = LinkValueCast(sv);
	69	if (lv)
	70	for (const Handle& g : lv->to_handle_seq())
	71	if (g->is_link() and g->get_arity() >= 3)
	72	out.push_back({g->getOutgoingAtom(0), g->getOutgoingAtom(2)});
	73	return out;
	74	}
	75
	76	// Execute a multi-variable MeetLink; return the set of FIRST-variable ($X)
	77	// groundings (results are ListLinks of the bound vars in order).
	78	static std::set<std::string> run_query_firstvar(AtomSpace* as, Handle vardecl, Handle body)
	79	{
	80	Handle meet = as->add_link(MEET_LINK, vardecl, body);
	81	ValuePtr sv = meet->execute(as);
	82	std::set<std::string> names;
	83	LinkValuePtr lv = LinkValueCast(sv);
	84	if (lv)
	85	for (const Handle& g : lv->to_handle_seq())
	86	{
	87	if (g->is_link() and g->get_arity() >= 1)
	88	names.insert(g->getOutgoingAtom(0)->get_name());
	89	else if (g->is_node())
	90	names.insert(g->get_name());
	91	}
	92	return names;
	93	}
	94
	95	// Execute a MeetLink (vardecl, body) in `as`; return grounded node names.
	96	static std::set<std::string> run_query(AtomSpace* as, Handle vardecl, Handle body)
	97	{
	98	Handle meet = as->add_link(MEET_LINK, vardecl, body);
	99	ValuePtr sv = meet->execute(as);
	100	std::set<std::string> names;
	101	LinkValuePtr lv = LinkValueCast(sv);
	102	if (lv)
	103	for (const Handle& h : lv->to_handle_seq())
	104	if (h->is_node()) names.insert(h->get_name());
	105	return names;
	106	}
	107
	108	// Run (Meet $X:Concept (Inheritance $X animal)) in `as`; return the set
	109	// of grounded node-names. Used to compare RAM-baseline vs GPU-fetched data.
	110	static std::set<std::string> run_meet(AtomSpace* as, Handle animal)
	111	{
	112	Handle var = as->add_node(VARIABLE_NODE, "$X");
	113	Handle tv = as->add_link(TYPED_VARIABLE_LINK, var,
	114	as->add_node(TYPE_NODE, "ConceptNode"));
	115	Handle body = as->add_link(INHERITANCE_LINK, var, animal);
	116	Handle meet = as->add_link(MEET_LINK, tv, body);
	117	ValuePtr sv = meet->execute(as);
	118	std::set<std::string> names;
	119	LinkValuePtr lv = LinkValueCast(sv);
	120	if (lv)
	121	for (const Handle& h : lv->to_handle_seq())
	122	if (h->is_node()) names.insert(h->get_name());
	123	return names;
	124	}
	125
54	126	// ============================================================
55	127	// Test that open/close/connected work correctly.
56	128
…	…	@@ -279,4 +351,362 @@ void GpuBasicSaveUTest::test_roundtrip(void)
279	351	logger().debug("END TEST: %s", __FUNCTION__);
280	352	}
281	353
	354	// ============================================================
	355	// Phase 2b: fetch incoming-set from the GPU into a fresh AtomSpace.
	356
	357	void GpuBasicSaveUTest::test_fetch_incoming(void)
	358	{
	359	logger().debug("BEGIN TEST: %s", __FUNCTION__);
	360
	361	GpuStorageNode* store = new GpuStorageNode(GPU_STORAGE_NODE, "gpu://test");
	362	store->open();
	363
	364	// Knowledge graph: cat is-a animal, dog is-a animal, salmon is-a fish.
	365	AtomSpace* as1 = new AtomSpace();
	366	Handle cat = as1->add_node(CONCEPT_NODE, "cat");
	367	Handle dog = as1->add_node(CONCEPT_NODE, "dog");
	368	Handle animal = as1->add_node(CONCEPT_NODE, "animal");
	369	Handle fish = as1->add_node(CONCEPT_NODE, "fish");
	370	Handle salmon = as1->add_node(CONCEPT_NODE, "salmon");
	371	as1->add_link(INHERITANCE_LINK, cat, animal);
	372	as1->add_link(INHERITANCE_LINK, dog, animal);
	373	as1->add_link(INHERITANCE_LINK, salmon, fish);
	374
	375	store->storeAtomSpace(as1);
	376	store->barrier();
	377
	378	// Fresh AtomSpace that only knows 'animal'. Pull its incoming
	379	// InheritanceLinks back from the GPU.
	380	AtomSpace* as2 = new AtomSpace();
	381	Handle animal2 = as2->add_node(CONCEPT_NODE, "animal");
	382	store->fetchIncomingByType(as2, animal2, INHERITANCE_LINK);
	383
	384	Handle inh_cat = as2->get_atom(createLink(INHERITANCE_LINK,
	385	createNode(CONCEPT_NODE, "cat"), createNode(CONCEPT_NODE, "animal")));
	386	Handle inh_dog = as2->get_atom(createLink(INHERITANCE_LINK,
	387	createNode(CONCEPT_NODE, "dog"), createNode(CONCEPT_NODE, "animal")));
	388	Handle inh_salmon = as2->get_atom(createLink(INHERITANCE_LINK,
	389	createNode(CONCEPT_NODE, "salmon"), createNode(CONCEPT_NODE, "fish")));
	390
	391	TSM_ASSERT("cat is-a animal should be fetched", inh_cat != nullptr);
	392	TSM_ASSERT("dog is-a animal should be fetched", inh_dog != nullptr);
	393	TSM_ASSERT("salmon is-a fish must NOT be fetched (not in animal's incoming)",
	394	inh_salmon == nullptr);
	395
	396	// animal's incoming-set has exactly 2 InheritanceLinks.
	397	TSM_ASSERT_EQUALS("animal incoming-by-type count",
	398	animal2->getIncomingSetSizeByType(INHERITANCE_LINK), (size_t)2);
	399
	400	store->close();
	401	delete store;
	402	delete as1;
	403	delete as2;
	404
	405	printf("test_fetch_incoming: PASSED\n");
	406	logger().debug("END TEST: %s", __FUNCTION__);
	407	}
	408
	409	// ============================================================
	410	// Phase 2c: GOLD CHECK. Run the SAME query (a) in RAM (baseline) and
	411	// (b) against data fetched from the GPU. Results must be identical.
	412
	413	void GpuBasicSaveUTest::test_query(void)
	414	{
	415	logger().debug("BEGIN TEST: %s", __FUNCTION__);
	416
	417	GpuStorageNode* store = new GpuStorageNode(GPU_STORAGE_NODE, "gpu://test");
	418	store->open();
	419
	420	// KG: cat/dog/bear is-a animal; salmon is-a fish.
	421	// Use createAtomSpace() (shared_ptr) — MeetLink::execute needs shared_from_this.
	422	AtomSpacePtr as1 = createAtomSpace();
	423	Handle cat = as1->add_node(CONCEPT_NODE, "cat");
	424	Handle dog = as1->add_node(CONCEPT_NODE, "dog");
	425	Handle bear = as1->add_node(CONCEPT_NODE, "bear");
	426	Handle animal = as1->add_node(CONCEPT_NODE, "animal");
	427	Handle fish = as1->add_node(CONCEPT_NODE, "fish");
	428	Handle salmon = as1->add_node(CONCEPT_NODE, "salmon");
	429	as1->add_link(INHERITANCE_LINK, cat, animal);
	430	as1->add_link(INHERITANCE_LINK, dog, animal);
	431	as1->add_link(INHERITANCE_LINK, bear, animal);
	432	as1->add_link(INHERITANCE_LINK, salmon, fish);
	433
	434	store->storeAtomSpace(as1.get());
	435	store->barrier();
	436
	437	// (a) Baseline: query the full AtomSpace in RAM.
	438	std::set<std::string> baseline = run_meet(as1.get(), animal);
	439
	440	// (b) GPU path: fresh AtomSpace, pull animal's incoming from the GPU,
	441	// then run the identical query against the fetched data.
	442	AtomSpacePtr as2 = createAtomSpace();
	443	Handle animal2 = as2->add_node(CONCEPT_NODE, "animal");
	444	store->fetchIncomingByType(as2.get(), animal2, INHERITANCE_LINK);
	445	std::set<std::string> gpu = run_meet(as2.get(), animal2);
	446
	447	std::set<std::string> expect = {"cat", "dog", "bear"};
	448	printf("test_query: baseline={");
	449	for (auto& s : baseline) printf("%s ", s.c_str());
	450	printf("} gpu={");
	451	for (auto& s : gpu) printf("%s ", s.c_str());
	452	printf("}\n");
	453
	454	TSM_ASSERT_EQUALS("baseline size should be 3", baseline.size(), (size_t)3);
	455	TSM_ASSERT("baseline == {cat,dog,bear}", baseline == expect);
	456	TSM_ASSERT("GOLD: GPU-path query result == RAM baseline", gpu == baseline);
	457	TSM_ASSERT("GPU-path == {cat,dog,bear}", gpu == expect);
	458
	459	store->close();
	460	delete store;
	461
	462	printf("test_query: PASSED (gold check: GPU query == in-RAM baseline)\n");
	463	logger().debug("END TEST: %s", __FUNCTION__);
	464	}
	465
	466	// ============================================================
	467	// V1 VALUE TEST: crisp multi-hop reasoning, no hallucination.
	468	// A multi-relation KG round-tripped through the GPU; multi-clause
	469	// queries with deliberate LLM-traps. Exact + deterministic.
	470
	471	void GpuBasicSaveUTest::test_v1_multihop(void)
	472	{
	473	logger().debug("BEGIN TEST: %s", __FUNCTION__);
	474
	475	GpuStorageNode* store = new GpuStorageNode(GPU_STORAGE_NODE, "gpu://test");
	476	store->open();
	477
	478	AtomSpacePtr as1 = createAtomSpace();
	479	// is-a (InheritanceLink)
	480	auto isa = [&](const char* x, const char* t) {
	481	as1->add_link(INHERITANCE_LINK,
	482	as1->add_node(CONCEPT_NODE, x), as1->add_node(CONCEPT_NODE, t)); };
	483	// lives-in (EvaluationLink (Predicate "lives-in") (List x place))
	484	auto lives = [&](const char* x, const char* p) {
	485	as1->add_link(EVALUATION_LINK, as1->add_node(PREDICATE_NODE, "lives-in"),
	486	as1->add_link(LIST_LINK,
	487	as1->add_node(CONCEPT_NODE, x), as1->add_node(CONCEPT_NODE, p))); };
	488
	489	isa("whale","mammal"); isa("dolphin","mammal"); isa("seal","mammal");
	490	isa("bat","mammal"); isa("cat","mammal");
	491	isa("shark","fish"); isa("salmon","fish"); isa("eagle","bird");
	492	lives("whale","water"); lives("dolphin","water"); lives("seal","water");
	493	lives("shark","water"); lives("salmon","water");
	494	lives("bat","air"); lives("eagle","air"); lives("cat","land");
	495
	496	// Store the whole KG to the GPU, then load it back from the GPU.
	497	store->storeAtomSpace(as1.get());
	498	store->barrier();
	499	AtomSpacePtr gas = createAtomSpace(); // GPU-sourced AtomSpace
	500	store->loadAtomSpace(gas.get());
	501
	502	// Multi-hop query builder: (is-a $X type) AND (lives-in $X place)
	503	auto query = [](AtomSpace* as, const char* type, const char* place) {
	504	Handle X = as->add_node(VARIABLE_NODE, "$X");
	505	Handle vd = as->add_link(TYPED_VARIABLE_LINK, X,
	506	as->add_node(TYPE_NODE, "ConceptNode"));
	507	Handle c1 = as->add_link(INHERITANCE_LINK, X, as->add_node(CONCEPT_NODE, type));
	508	Handle c2 = as->add_link(EVALUATION_LINK, as->add_node(PREDICATE_NODE, "lives-in"),
	509	as->add_link(LIST_LINK, X, as->add_node(CONCEPT_NODE, place)));
	510	Handle body = as->add_link(AND_LINK, c1, c2);
	511	return run_query(as, vd, body);
	512	};
	513
	514	// Q1: mammals that live in water. Traps: shark/salmon (water, but fish),
	515	// bat/cat (mammal, but not water).
	516	std::set<std::string> q1_base = query(as1.get(), "mammal", "water");
	517	std::set<std::string> q1_gpu = query(gas.get(), "mammal", "water");
	518	std::set<std::string> q1_exp = {"whale", "dolphin", "seal"};
	519
	520	printf("\nV1 Q1 'mammals that live in water': {");
	521	for (auto& s : q1_gpu) printf("%s ", s.c_str());
	522	printf("} (LLM trap: shark/salmon live in water but are fish)\n");
	523
	524	TSM_ASSERT("Q1 GPU == in-RAM baseline", q1_gpu == q1_base);
	525	TSM_ASSERT("Q1 == {whale,dolphin,seal}", q1_gpu == q1_exp);
	526
	527	// Q2: fish that live in water. Trap: whale/dolphin (water, but mammals).
	528	std::set<std::string> q2_gpu = query(gas.get(), "fish", "water");
	529	std::set<std::string> q2_exp = {"shark", "salmon"};
	530	printf("V1 Q2 'fish that live in water': {");
	531	for (auto& s : q2_gpu) printf("%s ", s.c_str());
	532	printf("} (LLM trap: whale/dolphin live in water but are mammals)\n");
	533	TSM_ASSERT("Q2 == {shark,salmon}", q2_gpu == q2_exp);
	534
	535	// Determinism: same query, run 5x, identical every time.
	536	bool stable = true;
	537	for (int i = 0; i < 5; i++)
	538	if (query(gas.get(), "mammal", "water") != q1_exp) stable = false;
	539	TSM_ASSERT("Q1 deterministic over 5 runs", stable);
	540	printf("V1 determinism: 5/5 identical runs -> %s\n", stable ? "PASS" : "FAIL");
	541
	542	store->close();
	543	delete store;
	544
	545	printf("test_v1_multihop: PASSED (crisp multi-hop, grounded, deterministic)\n");
	546	logger().debug("END TEST: %s", __FUNCTION__);
	547	}
	548
	549	// ============================================================
	550	// V4 NEUROSYMBOLIC LOOP: facts extracted by an LLM from text are stored
	551	// in the GPU AtomSpace; a 3-fact composition query (which an LLM-alone
	552	// is prone to get wrong) is answered exactly by the symbolic engine.
	553	//
	554	// Source text (the LLM extraction is done outside this test):
	555	// "Alice manages Bob, who manages Carol; Carol in turn manages Dave.
	556	// Alice and Bob are based in the Boston office, while Carol and Dave
	557	// work out of Seattle."
	558	// Question: who manages someone in Seattle but is themselves in Boston?
	559	// = (manages $X $Y) AND (works-in $Y Seattle) AND (works-in $X Boston)
	560	// Answer: Bob (manages Carol@Seattle; Bob@Boston). Needs 3 facts about
	561	// TWO different people — the cross-referencing LLMs bungle.
	562
	563	void GpuBasicSaveUTest::test_v4_neurosymbolic(void)
	564	{
	565	logger().debug("BEGIN TEST: %s", __FUNCTION__);
	566
	567	GpuStorageNode* store = new GpuStorageNode(GPU_STORAGE_NODE, "gpu://test");
	568	store->open();
	569
	570	AtomSpacePtr as1 = createAtomSpace();
	571	auto rel = [&](const char* pred, const char* a, const char* b) {
	572	as1->add_link(EVALUATION_LINK, as1->add_node(PREDICATE_NODE, pred),
	573	as1->add_link(LIST_LINK,
	574	as1->add_node(CONCEPT_NODE, a), as1->add_node(CONCEPT_NODE, b))); };
	575
	576	// LLM-extracted facts:
	577	rel("manages","Alice","Bob"); rel("manages","Bob","Carol"); rel("manages","Carol","Dave");
	578	rel("works-in","Alice","Boston"); rel("works-in","Bob","Boston");
	579	rel("works-in","Carol","Seattle"); rel("works-in","Dave","Seattle");
	580
	581	// Round-trip through the GPU.
	582	store->storeAtomSpace(as1.get());
	583	store->barrier();
	584	AtomSpacePtr gas = createAtomSpace();
	585	store->loadAtomSpace(gas.get());
	586
	587	auto query = [](AtomSpace* as) {
	588	Handle X = as->add_node(VARIABLE_NODE, "$X");
	589	Handle Y = as->add_node(VARIABLE_NODE, "$Y");
	590	Handle vd = as->add_link(VARIABLE_LIST,
	591	as->add_link(TYPED_VARIABLE_LINK, X, as->add_node(TYPE_NODE, "ConceptNode")),
	592	as->add_link(TYPED_VARIABLE_LINK, Y, as->add_node(TYPE_NODE, "ConceptNode")));
	593	Handle mgr = as->add_node(PREDICATE_NODE, "manages");
	594	Handle win = as->add_node(PREDICATE_NODE, "works-in");
	595	Handle c1 = as->add_link(EVALUATION_LINK, mgr, as->add_link(LIST_LINK, X, Y));
	596	Handle c2 = as->add_link(EVALUATION_LINK, win,
	597	as->add_link(LIST_LINK, Y, as->add_node(CONCEPT_NODE, "Seattle")));
	598	Handle c3 = as->add_link(EVALUATION_LINK, win,
	599	as->add_link(LIST_LINK, X, as->add_node(CONCEPT_NODE, "Boston")));
	600	Handle body = as->add_link(AND_LINK, c1, c2, c3);
	601	return run_query_firstvar(as, vd, body);
	602	};
	603
	604	std::set<std::string> ans = query(gas.get());
	605	std::set<std::string> exp = {"Bob"};
	606
	607	printf("\nV4 neurosymbolic: 'who manages someone in Seattle but works in Boston?'\n");
	608	printf(" AtomSpace answer: {");
	609	for (auto& s : ans) printf("%s ", s.c_str());
	610	printf("} (exact 3-fact composition; LLM-alone often says Alice or Carol)\n");
	611
	612	TSM_ASSERT("V4 answer == {Bob}", ans == exp);
	613	TSM_ASSERT("V4 GPU == baseline", ans == query(as1.get()));
	614
	615	store->close();
	616	delete store;
	617
	618	printf("test_v4_neurosymbolic: PASSED (text->facts->GPU AtomSpace->exact 3-hop answer)\n");
	619	logger().debug("END TEST: %s", __FUNCTION__);
	620	}
	621
	622	// ============================================================
	623	// Phase 3 AGENT LOOP: a minimal Linas 5-bullet agent over a GPU world-model.
	624	// world-model = GpuStorageNode AtomSpace (inside) ; sense = ingest facts ;
	625	// decide = query the world-model ; act = derive a new fact ;
	626	// update = store it back to the GPU. Loop to fixpoint.
	627	// Task: forward-chain the transitive closure of 'manages' (A->B->C->D).
	628
	629	void GpuBasicSaveUTest::test_agent_loop(void)
	630	{
	631	logger().debug("BEGIN TEST: %s", __FUNCTION__);
	632
	633	GpuStorageNode* store = new GpuStorageNode(GPU_STORAGE_NODE, "gpu://test");
	634	store->open();
	635
	636	AtomSpacePtr as = createAtomSpace(); // the agent's world-model
	637	Handle mgr = as->add_node(PREDICATE_NODE, "manages");
	638	Handle rch = as->add_node(PREDICATE_NODE, "reaches");
	639	auto mk = [&](Handle pred, const char* a, const char* b) {
	640	return as->add_link(EVALUATION_LINK, pred,
	641	as->add_link(LIST_LINK,
	642	as->add_node(CONCEPT_NODE, a), as->add_node(CONCEPT_NODE, b))); };
	643
	644	// SENSE: a chain of management facts. Base case: reaches = manages.
	645	const char* chain[3][2] = {{"A","B"},{"B","C"},{"C","D"}};
	646	std::set<std::pair<std::string,std::string>> known;
	647	for (auto& e : chain) { mk(mgr, e[0], e[1]); mk(rch, e[0], e[1]);
	648	known.insert({e[0], e[1]}); }
	649	store->storeAtomSpace(as.get());
	650	store->barrier();
	651
	652	// DECIDE pattern: reaches($X,$Y) AND manages($Y,$Z) -> derive reaches($X,$Z)
	653	Handle X = as->add_node(VARIABLE_NODE, "$X");
	654	Handle Y = as->add_node(VARIABLE_NODE, "$Y");
	655	Handle Z = as->add_node(VARIABLE_NODE, "$Z");
	656	Handle ct = as->add_node(TYPE_NODE, "ConceptNode");
	657	Handle vd = as->add_link(VARIABLE_LIST,
	658	as->add_link(TYPED_VARIABLE_LINK, X, ct),
	659	as->add_link(TYPED_VARIABLE_LINK, Y, ct),
	660	as->add_link(TYPED_VARIABLE_LINK, Z, ct));
	661	Handle body = as->add_link(AND_LINK,
	662	as->add_link(EVALUATION_LINK, rch, as->add_link(LIST_LINK, X, Y)),
	663	as->add_link(EVALUATION_LINK, mgr, as->add_link(LIST_LINK, Y, Z)));
	664
	665	int cycles = 0, total = 0;
	666	while (cycles <= 10)
	667	{
	668	auto xz = run_xz(as.get(), vd, body); // DECIDE (query)
	669	int added = 0;
	670	for (auto& p : xz)
	671	{
	672	std::string xn = p.first->get_name(), zn = p.second->get_name();
	673	if (known.count({xn, zn})) continue;
	674	Handle nf = as->add_link(EVALUATION_LINK, rch, // ACT (derive)
	675	as->add_link(LIST_LINK, p.first, p.second));
	676	store->storeAtom(nf, true); // UPDATE (persist to GPU)
	677	known.insert({xn, zn});
	678	added++;
	679	}
	680	cycles++; total += added;
	681	if (0 == added) break;
	682	}
	683	store->barrier();
	684
	685	printf("\nPhase 3 agent loop: %d cycles, derived %d new 'reaches' facts\n", cycles, total);
	686	printf(" transitive closure: ");
	687	for (auto& k : known) printf("%s->%s ", k.first.c_str(), k.second.c_str());
	688	printf("\n");
	689
	690	// Transitive closure of A->B->C->D has 6 pairs; A->D is the 3-hop derivation.
	691	TSM_ASSERT_EQUALS("closure size == 6", known.size(), (size_t)6);
	692	TSM_ASSERT("derived A->D (3-hop)", known.count({"A","D"}) == 1);
	693
	694	// Persistence: derived facts live on the GPU. Load fresh, count 'reaches'.
	695	AtomSpacePtr fresh = createAtomSpace();
	696	store->loadAtomSpace(fresh.get());
	697	HandleSeq evals;
	698	fresh->get_handles_by_type(evals, EVALUATION_LINK, false);
	699	int rcount = 0;
	700	for (auto& h : evals)
	701	if (h->getOutgoingAtom(0)->get_name() == "reaches") rcount++;
	702	printf(" 'reaches' facts persisted on GPU: %d\n", rcount);
	703	TSM_ASSERT_EQUALS("6 reaches persisted on GPU", rcount, 6);
	704
	705	store->close();
	706	delete store;
	707
	708	printf("test_agent_loop: PASSED (sense->query->derive->update closure on GPU world-model)\n");
	709	logger().debug("END TEST: %s", __FUNCTION__);
	710	}
	711
282	712	/* ============================= END OF FILE ================= */

tests/persist/gpu/test_incoming.cu +79 -0

…	…	@@ -0,0 +1,79 @@
	1	/*
	2	* test_incoming.cu — standalone test for GpuAtomTable incoming-set index (Phase 2a).
	3	*
	4	* Slots: a=0 b=1 c=2 (nodes);
	5	* L1=3 (a,b) L2=4 (b,c) L3=5 (a,b,c) L4=6 (L1,c) [link-of-link]
	6	* Expected incoming sets (link-slots referencing each atom):
	7	* a:{3,5} b:{3,4,5} c:{4,5,6} L1:{6} L2:{} L3:{} L4:{}
	8	*
	9	* Build: nvcc -DHAVE_CUDA -I../../.. test_incoming.cu \
	10	* ../../../opencog/persist/gpu/GpuAtomTable.cu -arch=sm_75 -o test_incoming
	11	*/
	12	#include "opencog/persist/gpu/GpuAtomTable.h"
	13	#include <cstdio>
	14	#include <vector>
	15	#include <algorithm>
	16
	17	static int fails = 0;
	18	static void check_set(const char* name, GpuAtomTable* t, uint32_t slot,
	19	std::vector<uint32_t> expect)
	20	{
	21	uint32_t buf[64]; uint32_t count = 0;
	22	cuda_fetch_incoming(t, slot, buf, 64, &count);
	23	std::vector<uint32_t> got(buf, buf + (count < 64 ? count : 64));
	24	std::sort(got.begin(), got.end());
	25	std::sort(expect.begin(), expect.end());
	26	if (got == expect) { printf("PASS incoming(%s) = {", name);
	27	for (auto x : got) printf("%u ", x); printf("}\n"); }
	28	else { fails++; printf("FAIL incoming(%s): got {", name);
	29	for (auto x : got) printf("%u ", x); printf("} expected {");
	30	for (auto x : expect) printf("%u ", x); printf("}\n"); }
	31	}
	32
	33	int main()
	34	{
	35	printf("=== Phase 2a: GpuAtomTable incoming-set index ===\n\n");
	36	GpuAtomTable t;
	37	if (cuda_table_alloc(&t) != 0) { printf("alloc failed\n"); return 1; }
	38
	39	cuda_store_node(&t, 0, 1, "a", 1);
	40	cuda_store_node(&t, 1, 1, "b", 1);
	41	cuda_store_node(&t, 2, 1, "c", 1);
	42	uint32_t L1[] = {0,1}; cuda_store_link(&t, 3, 2, L1, 2);
	43	uint32_t L2[] = {1,2}; cuda_store_link(&t, 4, 2, L2, 2);
	44	uint32_t L3[] = {0,1,2}; cuda_store_link(&t, 5, 2, L3, 3);
	45	uint32_t L4[] = {3,2}; cuda_store_link(&t, 6, 2, L4, 2); // contains link L1
	46
	47	if (cuda_table_build_incoming(&t) != 0) { printf("build_incoming failed\n"); return 1; }
	48	printf("atom_count=%u in_pool_used=%u in_built=%u\n\n",
	49	t.atom_count, t.in_pool_used, t.in_built);
	50
	51	check_set("a", &t, 0, {3,5});
	52	check_set("b", &t, 1, {3,4,5});
	53	check_set("c", &t, 2, {4,5,6});
	54	check_set("L1", &t, 3, {6});
	55	check_set("L2", &t, 4, {});
	56	check_set("L3", &t, 5, {});
	57	check_set("L4", &t, 6, {});
	58
	59	// GPU-resident JOIN: intersect incoming sets entirely on the GPU.
	60	{
	61	uint32_t buf[64]; uint32_t cnt;
	62	auto chk = [&](const char* nm, uint32_t a, uint32_t b, std::vector<uint32_t> exp) {
	63	cnt = 0; cuda_intersect_incoming(&t, a, b, buf, 64, &cnt);
	64	std::vector<uint32_t> got(buf, buf + (cnt < 64 ? cnt : 64));
	65	std::sort(got.begin(), got.end()); std::sort(exp.begin(), exp.end());
	66	if (got == exp) { printf("PASS intersect(%s) = {", nm);
	67	for (auto x : got) printf("%u ", x); printf("}\n"); }
	68	else { fails++; printf("FAIL intersect(%s)\n", nm); }
	69	};
	70	chk("a,b", 0, 1, {3,5}); // links containing both a and b: L1,L3
	71	chk("b,c", 1, 2, {4,5}); // L2,L3
	72	chk("a,c", 0, 2, {5}); // only L3 contains both a and c
	73	}
	74
	75	cuda_table_free(&t);
	76	printf("\n%s\n", fails ? "SOME CHECKS FAILED" :
	77	"ALL CHECKS PASSED - incoming-set index works on the real GpuAtomTable.");
	78	return fails ? 1 : 0;
	79	}

atomspace-gpu commits / fecc49c

Changed files (8) +859 -2