/*
* opencog/persist/gpu/GpuAtomTable.cu
*
* CUDA implementation of the GpuAtomTable API.
*
* AoS for per-atom fixed fields: one GpuAtom struct per slot,
* one cudaMemcpy to store/fetch all fields of one atom.
* Variable-length data (names, outgoing sets) in side pools.
*
* Copyright (C) 2026 OpenCog Foundation
*
* SPDX-License-Identifier: AGPL-3.0-or-later
*/
#include "GpuAtomTable.h"
#include <cuda_runtime.h>
#include <cstdio>
#include <cstring>
#include <vector>
/* ---------------------------------------------------------------- */
#define CUDA_CHECK(call, ctx) do { \
cudaError_t _e = (call); \
if (_e != cudaSuccess) { \
fprintf(stderr, "CUDA error in %s: %s\n", \
(ctx), cudaGetErrorString(_e)); \
return -1; \
} \
} while(0)
#define CUDA_CHECK_VOID(call, ctx) do { \
cudaError_t _e = (call); \
if (_e != cudaSuccess) { \
fprintf(stderr, "CUDA error in %s: %s\n", \
(ctx), cudaGetErrorString(_e)); \
} \
} while(0)
/* ================================================================
* gpu_table_alloc
* ================================================================ */
extern "C"
int cuda_table_alloc(GpuAtomTable* t)
{
memset(t, 0, sizeof(GpuAtomTable));
CUDA_CHECK(cudaMalloc(&t->atoms, GPU_ATOM_CAPACITY * sizeof(GpuAtom)),
"alloc atoms");
CUDA_CHECK(cudaMalloc(&t->name_pool, GPU_NAME_POOL_BYTES),
"alloc name_pool");
CUDA_CHECK(cudaMalloc(&t->out_pool, GPU_OUT_POOL_SLOTS * sizeof(uint32_t)),
"alloc out_pool");
CUDA_CHECK(cudaMalloc(&t->in_off, (GPU_ATOM_CAPACITY + 1) * sizeof(uint32_t)),
"alloc in_off");
CUDA_CHECK(cudaMalloc(&t->in_pool, GPU_IN_POOL_SLOTS * sizeof(uint32_t)),
"alloc in_pool");
CUDA_CHECK(cudaMemset(t->atoms, 0, GPU_ATOM_CAPACITY * sizeof(GpuAtom)),
"zero atoms");
CUDA_CHECK(cudaMemset(t->name_pool, 0, GPU_NAME_POOL_BYTES),
"zero name_pool");
CUDA_CHECK(cudaMemset(t->out_pool, 0, GPU_OUT_POOL_SLOTS * sizeof(uint32_t)),
"zero out_pool");
CUDA_CHECK(cudaMemset(t->in_off, 0, (GPU_ATOM_CAPACITY + 1) * sizeof(uint32_t)),
"zero in_off");
t->atom_count = 0;
t->name_pool_used = 0;
t->out_pool_used = 0;
t->in_pool_used = 0;
t->in_built = 0;
return 0;
}
/* ================================================================
* gpu_table_free
* ================================================================ */
extern "C"
void cuda_table_free(GpuAtomTable* t)
{
if (t->atoms) cudaFree(t->atoms);
if (t->name_pool) cudaFree(t->name_pool);
if (t->out_pool) cudaFree(t->out_pool);
if (t->in_off) cudaFree(t->in_off);
if (t->in_pool) cudaFree(t->in_pool);
memset(t, 0, sizeof(GpuAtomTable));
}
/* ================================================================
* gpu_table_clear
* ================================================================ */
extern "C"
void cuda_table_clear(GpuAtomTable* t)
{
CUDA_CHECK_VOID(cudaMemset(t->atoms, 0, GPU_ATOM_CAPACITY * sizeof(GpuAtom)),
"clear atoms");
CUDA_CHECK_VOID(cudaMemset(t->name_pool, 0, GPU_NAME_POOL_BYTES),
"clear name_pool");
CUDA_CHECK_VOID(cudaMemset(t->out_pool, 0, GPU_OUT_POOL_SLOTS * sizeof(uint32_t)),
"clear out_pool");
CUDA_CHECK_VOID(cudaMemset(t->in_off, 0, (GPU_ATOM_CAPACITY + 1) * sizeof(uint32_t)),
"clear in_off");
t->atom_count = 0;
t->name_pool_used = 0;
t->out_pool_used = 0;
t->in_pool_used = 0;
t->in_built = 0;
}
/* ================================================================
* gpu_store_node -- store type + name at slot.
* ================================================================ */
extern "C"
int cuda_store_node(GpuAtomTable* t, uint32_t slot,
uint16_t type, const char* name, uint16_t name_len)
{
if (slot >= GPU_ATOM_CAPACITY) return -1;
if (t->name_pool_used + name_len > GPU_NAME_POOL_BYTES) return -1;
uint32_t off = t->name_pool_used;
/* Build the atom struct on host, copy in one shot */
GpuAtom a;
memset(&a, 0, sizeof(a));
a.type = type;
a.is_node = 1;
a.data_offset = off;
a.data_len = name_len;
CUDA_CHECK(cudaMemcpy(t->atoms + slot, &a, sizeof(GpuAtom),
cudaMemcpyHostToDevice), "store node atom");
/* Append name to pool */
if (name_len > 0)
{
CUDA_CHECK(cudaMemcpy(t->name_pool + off, name, name_len,
cudaMemcpyHostToDevice), "store node name");
}
t->name_pool_used += name_len;
if (slot >= t->atom_count) t->atom_count = slot + 1;
return 0;
}
/* ================================================================
* gpu_store_link -- store type + outgoing set at slot.
* ================================================================ */
extern "C"
int cuda_store_link(GpuAtomTable* t, uint32_t slot,
uint16_t type,
const uint32_t* outgoing, uint16_t arity)
{
if (slot >= GPU_ATOM_CAPACITY) return -1;
if (t->out_pool_used + arity > GPU_OUT_POOL_SLOTS) return -1;
uint32_t off = t->out_pool_used;
GpuAtom a;
memset(&a, 0, sizeof(a));
a.type = type;
a.is_node = 0;
a.data_offset = off;
a.data_len = arity;
CUDA_CHECK(cudaMemcpy(t->atoms + slot, &a, sizeof(GpuAtom),
cudaMemcpyHostToDevice), "store link atom");
/* Append outgoing set to pool */
if (arity > 0)
{
CUDA_CHECK(cudaMemcpy(t->out_pool + off, outgoing,
arity * sizeof(uint32_t),
cudaMemcpyHostToDevice), "store link outgoing");
}
t->out_pool_used += arity;
if (slot >= t->atom_count) t->atom_count = slot + 1;
return 0;
}
/* ================================================================
* gpu_fetch_node -- read node back to host.
* ================================================================ */
extern "C"
int cuda_fetch_node(const GpuAtomTable* t, uint32_t slot,
uint16_t* type, char* name_buf, uint16_t* name_len)
{
if (slot >= GPU_ATOM_CAPACITY) return -1;
GpuAtom a;
CUDA_CHECK(cudaMemcpy(&a, t->atoms + slot, sizeof(GpuAtom),
cudaMemcpyDeviceToHost), "fetch node atom");
*type = a.type;
*name_len = a.data_len;
if (a.data_len > 0 && name_buf)
{
CUDA_CHECK(cudaMemcpy(name_buf, t->name_pool + a.data_offset,
a.data_len,
cudaMemcpyDeviceToHost), "fetch node name");
}
return 0;
}
/* ================================================================
* gpu_fetch_link -- read link back to host.
* ================================================================ */
extern "C"
int cuda_fetch_link(const GpuAtomTable* t, uint32_t slot,
uint16_t* type, uint32_t* out_buf, uint16_t* arity)
{
if (slot >= GPU_ATOM_CAPACITY) return -1;
GpuAtom a;
CUDA_CHECK(cudaMemcpy(&a, t->atoms + slot, sizeof(GpuAtom),
cudaMemcpyDeviceToHost), "fetch link atom");
*type = a.type;
*arity = a.data_len;
if (a.data_len > 0 && out_buf)
{
CUDA_CHECK(cudaMemcpy(out_buf, t->out_pool + a.data_offset,
a.data_len * sizeof(uint32_t),
cudaMemcpyDeviceToHost), "fetch link outgoing");
}
return 0;
}
/* ================================================================
* gpu_table_barrier
* ================================================================ */
extern "C"
void cuda_table_barrier(GpuAtomTable* t)
{
(void)t;
cudaDeviceSynchronize();
}
/* ================================================================
* Incoming-set index (Phase 2a).
* Build a CSR of link-slots referencing each atom, from the stored
* outgoing sets. count -> exclusive-scan -> scatter.
* ================================================================ */
__global__ void k_count_incoming(const GpuAtom* atoms, const uint32_t* out_pool,
uint32_t n, uint32_t* counts)
{
uint32_t s = blockIdx.x * blockDim.x + threadIdx.x;
if (s >= n) return;
GpuAtom a = atoms[s];
if (a.is_node) return;
for (uint16_t i = 0; i < a.data_len; i++)
atomicAdd(&counts[out_pool[a.data_offset + i]], 1u);
}
__global__ void k_scatter_incoming(const GpuAtom* atoms, const uint32_t* out_pool,
uint32_t n, uint32_t* cursor, uint32_t* in_pool)
{
uint32_t s = blockIdx.x * blockDim.x + threadIdx.x;
if (s >= n) return;
GpuAtom a = atoms[s];
if (a.is_node) return;
for (uint16_t i = 0; i < a.data_len; i++) {
uint32_t o = out_pool[a.data_offset + i];
uint32_t pos = atomicAdd(&cursor[o], 1u);
in_pool[pos] = s;
}
}
extern "C"
int cuda_table_build_incoming(GpuAtomTable* t)
{
uint32_t n = t->atom_count;
t->in_built = 0;
t->in_pool_used = 0;
if (n == 0) { t->in_built = 1; return 0; }
uint32_t* d_counts;
CUDA_CHECK(cudaMalloc(&d_counts, n * sizeof(uint32_t)), "alloc counts");
CUDA_CHECK(cudaMemset(d_counts, 0, n * sizeof(uint32_t)), "zero counts");
int TB = 256;
int GB = (int)((n + TB - 1) / TB);
k_count_incoming<<<GB, TB>>>(t->atoms, t->out_pool, n, d_counts);
CUDA_CHECK(cudaDeviceSynchronize(), "count incoming");
std::vector<uint32_t> counts(n);
CUDA_CHECK(cudaMemcpy(counts.data(), d_counts, n * sizeof(uint32_t),
cudaMemcpyDeviceToHost), "copy counts");
std::vector<uint32_t> off(n + 1);
off[0] = 0;
for (uint32_t i = 0; i < n; i++) off[i + 1] = off[i] + counts[i];
uint32_t total = off[n];
if (total > GPU_IN_POOL_SLOTS) { cudaFree(d_counts); return -1; }
CUDA_CHECK(cudaMemcpy(t->in_off, off.data(), (n + 1) * sizeof(uint32_t),
cudaMemcpyHostToDevice), "upload in_off");
uint32_t* d_cursor;
CUDA_CHECK(cudaMalloc(&d_cursor, n * sizeof(uint32_t)), "alloc cursor");
CUDA_CHECK(cudaMemcpy(d_cursor, off.data(), n * sizeof(uint32_t),
cudaMemcpyHostToDevice), "init cursor");
k_scatter_incoming<<<GB, TB>>>(t->atoms, t->out_pool, n, d_cursor, t->in_pool);
CUDA_CHECK(cudaDeviceSynchronize(), "scatter incoming");
cudaFree(d_counts);
cudaFree(d_cursor);
t->in_pool_used = total;
t->in_built = 1;
return 0;
}
extern "C"
int cuda_fetch_incoming(const GpuAtomTable* t, uint32_t slot,
uint32_t* buf, uint32_t max, uint32_t* count)
{
*count = 0;
if (slot >= t->atom_count) return -1;
uint32_t bounds[2];
CUDA_CHECK(cudaMemcpy(bounds, t->in_off + slot, 2 * sizeof(uint32_t),
cudaMemcpyDeviceToHost), "fetch in_off");
uint32_t c = bounds[1] - bounds[0];
*count = c;
uint32_t k = (c < max) ? c : max;
if (k > 0 && buf)
CUDA_CHECK(cudaMemcpy(buf, t->in_pool + bounds[0], k * sizeof(uint32_t),
cudaMemcpyDeviceToHost), "fetch incoming");
return 0;
}
/* GPU-resident join: intersect incoming(a) and incoming(b) on the GPU. */
__global__ void k_intersect_incoming(const uint32_t* in_off, const uint32_t* in_pool,
uint32_t a, uint32_t b,
uint32_t* out, int maxn, int* cnt)
{
uint32_t as = in_off[a], ae = in_off[a + 1];
uint32_t bs = in_off[b], be = in_off[b + 1];
for (uint32_t i = as + threadIdx.x; i < ae; i += blockDim.x) {
uint32_t x = in_pool[i];
bool found = false;
for (uint32_t j = bs; j < be; j++) if (in_pool[j] == x) { found = true; break; }
if (found) { int k = atomicAdd(cnt, 1); if (k < maxn) out[k] = x; }
}
}
extern "C"
int cuda_intersect_incoming(const GpuAtomTable* t, uint32_t a, uint32_t b,
uint32_t* outbuf, uint32_t max, uint32_t* count)
{
*count = 0;
if (a >= t->atom_count || b >= t->atom_count) return -1;
int* d_cnt;
uint32_t* d_out;
uint32_t alloc = (max > 0) ? max : 1;
CUDA_CHECK(cudaMalloc(&d_cnt, sizeof(int)), "alloc isect cnt");
CUDA_CHECK(cudaMemset(d_cnt, 0, sizeof(int)), "zero isect cnt");
CUDA_CHECK(cudaMalloc(&d_out, alloc * sizeof(uint32_t)), "alloc isect out");
k_intersect_incoming<<<1, 128>>>(t->in_off, t->in_pool, a, b, d_out, (int)max, d_cnt);
CUDA_CHECK(cudaDeviceSynchronize(), "intersect incoming");
int c;
CUDA_CHECK(cudaMemcpy(&c, d_cnt, sizeof(int), cudaMemcpyDeviceToHost), "copy isect cnt");
*count = (uint32_t)c;
uint32_t k = ((uint32_t)c < max) ? (uint32_t)c : max;
if (k > 0 && outbuf)
CUDA_CHECK(cudaMemcpy(outbuf, d_out, k * sizeof(uint32_t),
cudaMemcpyDeviceToHost), "copy isect out");
cudaFree(d_cnt);
cudaFree(d_out);
return 0;
}