/*
* gpu-atomspace-lib.c — Phase 7: GPU-Native AtomSpace Shared Library
*
* Wraps all 7 kernel files (Phases 0-6) into a single C API callable
* from Guile/Scheme via FFI. All atom pools and hash tables are
* GPU-resident; only small data transfers for input/output.
*
* Compile:
* gcc -shared -fPIC -o libgpu-atomspace.so gpu-atomspace-lib.c -lOpenCL -lm
*
* Test:
* gcc -DTEST_MAIN -o test-bridge gpu-atomspace-lib.c -lOpenCL -lm
*/
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <math.h>
/* ═══════════════════════════════════════════════════════════════════
* CAPACITY CONSTANTS (must match kernel -D flags)
* ═══════════════════════════════════════════════════════════════════ */
#define WORD_CAP 131072 /* 128K words */
#define PAIR_CAP 4194304 /* 4M pairs */
#define SEC_CAP 1048576 /* 1M sections */
#define WORD_HT_CAP 262144 /* 256K word HT */
#define PAIR_HT_CAP 8388608 /* 8M pair HT */
#define SEC_HT_CAP 2097152 /* 2M section HT */
#define DJH_HT_CAP 2097152 /* 2M disjunct HT */
#define CAND_CAP 524288 /* 512K candidates */
#define CAND_HT_CAP 1048576 /* 1M candidate HT */
#define CONN_WORD_CAP 2048 /* max words for connector cosine */
#define CONN_DIM_CAP 4096 /* max connector dimensions */
#define CONN_MAX_OUTPUT 100000 /* max output candidates */
#define MAX_BATCH_WORDS 200000
#define MAX_BATCH_SENTS 20000
#define MAX_BATCH_EDGES 200000
#define MAX_FILTER_OUT 65536
#define LOCAL_SIZE 256
/* ═══════════════════════════════════════════════════════════════════
* ERROR HANDLING
* ═══════════════════════════════════════════════════════════════════ */
#define CHK(err, msg) do { \
if ((err) != CL_SUCCESS) { \
fprintf(stderr, "GPU error: %s (code %d)\n", (msg), (err)); \
return -1; \
} \
} while(0)
static size_t round_up(size_t n) {
return ((n + LOCAL_SIZE - 1) / LOCAL_SIZE) * LOCAL_SIZE;
}
/* ═══════════════════════════════════════════════════════════════════
* GLOBAL STATE
* ═══════════════════════════════════════════════════════════════════ */
static cl_context g_ctx = NULL;
static cl_command_queue g_queue = NULL;
static cl_program g_prog = NULL;
static int g_init = 0;
/* ── Kernels ── */
static cl_kernel k_word_foc; /* word_find_or_create */
static cl_kernel k_count_sent; /* count_sentence_pairs */
static cl_kernel k_mi_res; /* compute_mi_resident */
static cl_kernel k_mi_dirty; /* compute_mi_dirty */
static cl_kernel k_mi_stats; /* mi_stats */
static cl_kernel k_extract_sec; /* extract_sections */
static cl_kernel k_word_norms; /* compute_word_norms */
static cl_kernel k_djh_chains; /* build_disjunct_chains */
static cl_kernel k_dot_prod; /* accumulate_dot_products */
static cl_kernel k_cosines; /* compute_cosines */
static cl_kernel k_filter_cand; /* filter_candidates */
static cl_kernel k_assign_cls; /* assign_classes */
static cl_kernel k_sub_pair; /* substitute_pairs */
static cl_kernel k_rebuild_idx; /* rebuild_pair_index */
static cl_kernel k_sub_sec; /* substitute_section_words */
static cl_kernel k_pool_stats; /* pool_stats */
static cl_kernel k_conn_norms; /* compute_connector_norms */
static cl_kernel k_conn_cosine; /* connector_cosine_brute */
/* ── GPU Buffers: WordPool ── */
static cl_mem b_wht_keys, b_wht_vals;
static cl_mem b_w_hash, b_w_count, b_w_class, b_w_free;
/* ── GPU Buffers: PairPool ── */
static cl_mem b_pht_keys, b_pht_vals;
static cl_mem b_p_wa, b_p_wb, b_p_count, b_p_mi, b_p_flags, b_p_free;
/* ── GPU Buffers: SectionPool ── */
static cl_mem b_sht_keys, b_sht_vals;
static cl_mem b_s_word, b_s_djh, b_s_count, b_s_free;
/* ── GPU Buffers: Cosine workspace ── */
static cl_mem b_norm_sq;
static cl_mem b_djh_keys, b_djh_vals, b_chain_next;
static cl_mem b_cht_keys, b_cht_vals;
static cl_mem b_c_wa, b_c_wb, b_c_dot, b_c_cos, b_c_free;
/* ── GPU Buffers: Counters ── */
static cl_mem b_total_pairs, b_total_secs;
static cl_mem b_num_changed, b_num_elim, b_num_merged;
/* ── GPU Buffers: Temp input ── */
static cl_mem b_flat_words, b_sent_off, b_sent_len;
static cl_mem b_in_hashes, b_out_indices;
static cl_mem b_edge_p1, b_edge_p2, b_edge_off, b_edge_cnt;
/* ── GPU Buffers: Filter output ── */
static cl_mem b_filt_wa, b_filt_wb, b_filt_cos, b_filt_cnt;
/* ── GPU Buffers: Stats output ── */
static cl_mem b_stats_out;
/* ── GPU Buffers: Connector-cosine workspace ── */
static cl_mem b_conn_matrix; /* float[CONN_WORD_CAP × CONN_DIM_CAP] */
static cl_mem b_conn_norms; /* float[CONN_WORD_CAP] */
static cl_mem b_cc_wa, b_cc_wb, b_cc_cos, b_cc_cnt; /* output candidates */
/* ═══════════════════════════════════════════════════════════════════
* FILE LOADING
* ═══════════════════════════════════════════════════════════════════ */
static char* load_file(const char* path) {
FILE* f = fopen(path, "r");
if (!f) { fprintf(stderr, "Cannot open %s\n", path); return NULL; }
fseek(f, 0, SEEK_END);
long len = ftell(f);
fseek(f, 0, SEEK_SET);
char* buf = (char*)malloc(len + 1);
fread(buf, 1, len, f);
buf[len] = '\0';
fclose(f);
return buf;
}
/* ═══════════════════════════════════════════════════════════════════
* BUFFER ALLOCATION HELPER
* ═══════════════════════════════════════════════════════════════════ */
static cl_mem mk_buf(size_t bytes, cl_int* err) {
return clCreateBuffer(g_ctx, CL_MEM_READ_WRITE, bytes, NULL, err);
}
/* ═══════════════════════════════════════════════════════════════════
* GPU INIT
* ═══════════════════════════════════════════════════════════════════ */
int gpu_init(const char* kernel_dir)
{
if (g_init) return 0; /* already initialized */
cl_int err;
cl_platform_id plat;
cl_device_id dev;
/* ── Platform & device ── */
err = clGetPlatformIDs(1, &plat, NULL);
CHK(err, "clGetPlatformIDs");
err = clGetDeviceIDs(plat, CL_DEVICE_TYPE_GPU, 1, &dev, NULL);
CHK(err, "clGetDeviceIDs");
char name[256];
clGetDeviceInfo(dev, CL_DEVICE_NAME, sizeof(name), name, NULL);
fprintf(stderr, "GPU: %s\n", name);
g_ctx = clCreateContext(NULL, 1, &dev, NULL, NULL, &err);
CHK(err, "clCreateContext");
g_queue = clCreateCommandQueue(g_ctx, dev, 0, &err);
CHK(err, "clCreateCommandQueue");
/* ── Load kernel sources ── */
const char* files[] = {
"gpu-hashtable.cl", "gpu-atomspace.cl", "gpu-counting.cl",
"gpu-mi.cl", "gpu-sections.cl", "gpu-cosine.cl", "gpu-substitute.cl",
"gpu-connector-cosine.cl"
};
int n_files = 8;
char** sources = (char**)malloc(sizeof(char*) * n_files);
size_t* lengths = (size_t*)malloc(sizeof(size_t) * n_files);
for (int i = 0; i < n_files; i++) {
char path[1024];
snprintf(path, sizeof(path), "%s/%s", kernel_dir, files[i]);
sources[i] = load_file(path);
if (!sources[i]) {
fprintf(stderr, "Failed to load %s\n", path);
return -1;
}
lengths[i] = strlen(sources[i]);
}
g_prog = clCreateProgramWithSource(g_ctx, n_files,
(const char**)sources, lengths, &err);
CHK(err, "clCreateProgramWithSource");
/* ── Build with capacity defines ── */
char opts[1024];
snprintf(opts, sizeof(opts),
"-DWORD_CAPACITY=%d -DPAIR_CAPACITY=%d -DSECTION_CAPACITY=%d "
"-DWORD_HT_CAPACITY=%d -DPAIR_HT_CAPACITY=%d -DSECTION_HT_CAPACITY=%d "
"-DDJH_HT_CAPACITY=%d -DCANDIDATE_CAPACITY=%d -DCANDIDATE_HT_CAPACITY=%d "
"-DMAX_CHAIN_LEN=%d -DMIN_NORM_SQ=50.0",
WORD_CAP, PAIR_CAP, SEC_CAP,
WORD_HT_CAP, PAIR_HT_CAP, SEC_HT_CAP,
DJH_HT_CAP, CAND_CAP, CAND_HT_CAP,
200);
err = clBuildProgram(g_prog, 1, &dev, opts, NULL, NULL);
if (err != CL_SUCCESS) {
char log[16384];
clGetProgramBuildInfo(g_prog, dev, CL_PROGRAM_BUILD_LOG,
sizeof(log), log, NULL);
fprintf(stderr, "Build error:\n%s\n", log);
return -1;
}
for (int i = 0; i < n_files; i++) free(sources[i]);
free(sources);
free(lengths);
fprintf(stderr, "Kernels compiled OK\n");
/* ── Create kernels ── */
#define MK_KERN(var, name) \
var = clCreateKernel(g_prog, name, &err); \
if (err != CL_SUCCESS) { \
fprintf(stderr, "Kernel '%s' failed: %d\n", name, err); \
return -1; \
}
MK_KERN(k_word_foc, "word_find_or_create");
MK_KERN(k_count_sent, "count_sentence_pairs");
MK_KERN(k_mi_res, "compute_mi_resident");
MK_KERN(k_mi_dirty, "compute_mi_dirty");
MK_KERN(k_mi_stats, "mi_stats");
MK_KERN(k_extract_sec, "extract_sections");
MK_KERN(k_word_norms, "compute_word_norms");
MK_KERN(k_djh_chains, "build_disjunct_chains");
MK_KERN(k_dot_prod, "accumulate_dot_products");
MK_KERN(k_cosines, "compute_cosines");
MK_KERN(k_filter_cand, "filter_candidates");
MK_KERN(k_assign_cls, "assign_classes");
MK_KERN(k_sub_pair, "substitute_pairs");
MK_KERN(k_rebuild_idx, "rebuild_pair_index");
MK_KERN(k_sub_sec, "substitute_section_words");
MK_KERN(k_pool_stats, "pool_stats");
MK_KERN(k_conn_norms, "compute_connector_norms");
MK_KERN(k_conn_cosine, "connector_cosine_brute");
fprintf(stderr, "18 kernels created\n");
/* ── Allocate GPU buffers ── */
#define BUF(var, bytes) \
var = mk_buf(bytes, &err); \
if (err != CL_SUCCESS) { \
fprintf(stderr, "Buffer alloc failed (%zu bytes): %d\n", \
(size_t)(bytes), err); \
return -1; \
}
/* WordPool */
BUF(b_wht_keys, sizeof(uint64_t) * WORD_HT_CAP);
BUF(b_wht_vals, sizeof(uint32_t) * WORD_HT_CAP);
BUF(b_w_hash, sizeof(uint64_t) * WORD_CAP);
BUF(b_w_count, sizeof(double) * WORD_CAP);
BUF(b_w_class, sizeof(uint32_t) * WORD_CAP);
BUF(b_w_free, sizeof(uint32_t));
/* PairPool */
BUF(b_pht_keys, sizeof(uint64_t) * PAIR_HT_CAP);
BUF(b_pht_vals, sizeof(uint32_t) * PAIR_HT_CAP);
BUF(b_p_wa, sizeof(uint32_t) * PAIR_CAP);
BUF(b_p_wb, sizeof(uint32_t) * PAIR_CAP);
BUF(b_p_count, sizeof(double) * PAIR_CAP);
BUF(b_p_mi, sizeof(double) * PAIR_CAP);
BUF(b_p_flags, sizeof(uint32_t) * PAIR_CAP);
BUF(b_p_free, sizeof(uint32_t));
/* SectionPool */
BUF(b_sht_keys, sizeof(uint64_t) * SEC_HT_CAP);
BUF(b_sht_vals, sizeof(uint32_t) * SEC_HT_CAP);
BUF(b_s_word, sizeof(uint32_t) * SEC_CAP);
BUF(b_s_djh, sizeof(uint64_t) * SEC_CAP);
BUF(b_s_count, sizeof(double) * SEC_CAP);
BUF(b_s_free, sizeof(uint32_t));
/* Cosine workspace */
BUF(b_norm_sq, sizeof(double) * WORD_CAP);
BUF(b_djh_keys, sizeof(uint64_t) * DJH_HT_CAP);
BUF(b_djh_vals, sizeof(uint32_t) * DJH_HT_CAP);
BUF(b_chain_next, sizeof(uint32_t) * SEC_CAP);
BUF(b_cht_keys, sizeof(uint64_t) * CAND_HT_CAP);
BUF(b_cht_vals, sizeof(uint32_t) * CAND_HT_CAP);
BUF(b_c_wa, sizeof(uint32_t) * CAND_CAP);
BUF(b_c_wb, sizeof(uint32_t) * CAND_CAP);
BUF(b_c_dot, sizeof(double) * CAND_CAP);
BUF(b_c_cos, sizeof(double) * CAND_CAP);
BUF(b_c_free, sizeof(uint32_t));
/* Counters */
BUF(b_total_pairs, sizeof(uint32_t));
BUF(b_total_secs, sizeof(uint32_t));
BUF(b_num_changed, sizeof(uint32_t));
BUF(b_num_elim, sizeof(uint32_t));
BUF(b_num_merged, sizeof(uint32_t));
/* Temp input */
BUF(b_flat_words, sizeof(uint32_t) * MAX_BATCH_WORDS);
BUF(b_sent_off, sizeof(uint32_t) * MAX_BATCH_SENTS);
BUF(b_sent_len, sizeof(uint32_t) * MAX_BATCH_SENTS);
BUF(b_in_hashes, sizeof(uint64_t) * MAX_BATCH_WORDS);
BUF(b_out_indices, sizeof(uint32_t) * MAX_BATCH_WORDS);
BUF(b_edge_p1, sizeof(uint32_t) * MAX_BATCH_EDGES);
BUF(b_edge_p2, sizeof(uint32_t) * MAX_BATCH_EDGES);
BUF(b_edge_off, sizeof(uint32_t) * MAX_BATCH_SENTS);
BUF(b_edge_cnt, sizeof(uint32_t) * MAX_BATCH_SENTS);
/* Filter output */
BUF(b_filt_wa, sizeof(uint32_t) * MAX_FILTER_OUT);
BUF(b_filt_wb, sizeof(uint32_t) * MAX_FILTER_OUT);
BUF(b_filt_cos, sizeof(double) * MAX_FILTER_OUT);
BUF(b_filt_cnt, sizeof(uint32_t));
/* Stats output */
BUF(b_stats_out, sizeof(uint32_t) * 4);
/* Connector-cosine workspace */
BUF(b_conn_matrix, sizeof(float) * CONN_WORD_CAP * CONN_DIM_CAP);
BUF(b_conn_norms, sizeof(float) * CONN_WORD_CAP);
BUF(b_cc_wa, sizeof(uint32_t) * CONN_MAX_OUTPUT);
BUF(b_cc_wb, sizeof(uint32_t) * CONN_MAX_OUTPUT);
BUF(b_cc_cos, sizeof(float) * CONN_MAX_OUTPUT);
BUF(b_cc_cnt, sizeof(uint32_t));
fprintf(stderr, "GPU buffers allocated\n");
/* ── Initialize buffers ── */
uint8_t ff = 0xFF;
uint32_t zero32 = 0;
double zero_d = 0.0;
/* Hash tables → 0xFF (empty sentinel) */
clEnqueueFillBuffer(g_queue, b_wht_keys, &ff, 1, 0,
sizeof(uint64_t) * WORD_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_wht_vals, &ff, 1, 0,
sizeof(uint32_t) * WORD_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_pht_keys, &ff, 1, 0,
sizeof(uint64_t) * PAIR_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_pht_vals, &ff, 1, 0,
sizeof(uint32_t) * PAIR_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_sht_keys, &ff, 1, 0,
sizeof(uint64_t) * SEC_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_sht_vals, &ff, 1, 0,
sizeof(uint32_t) * SEC_HT_CAP, 0, NULL, NULL);
/* Bump allocators → 0 */
clEnqueueFillBuffer(g_queue, b_w_free, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_p_free, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_s_free, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
/* Word counts and class IDs → 0 */
clEnqueueFillBuffer(g_queue, b_w_count, &zero_d, sizeof(double),
0, sizeof(double) * WORD_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_w_class, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t) * WORD_CAP, 0, NULL, NULL);
/* Counters → 0 */
clEnqueueFillBuffer(g_queue, b_total_pairs, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_total_secs, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clFinish(g_queue);
/* ── Set constant kernel args ── */
/* count_sentence_pairs: args 6-15 are constant pool buffers */
clSetKernelArg(k_count_sent, 6, sizeof(cl_mem), &b_pht_keys);
clSetKernelArg(k_count_sent, 7, sizeof(cl_mem), &b_pht_vals);
clSetKernelArg(k_count_sent, 8, sizeof(cl_mem), &b_p_wa);
clSetKernelArg(k_count_sent, 9, sizeof(cl_mem), &b_p_wb);
clSetKernelArg(k_count_sent, 10, sizeof(cl_mem), &b_p_count);
clSetKernelArg(k_count_sent, 11, sizeof(cl_mem), &b_p_mi);
clSetKernelArg(k_count_sent, 12, sizeof(cl_mem), &b_p_flags);
clSetKernelArg(k_count_sent, 13, sizeof(cl_mem), &b_p_free);
clSetKernelArg(k_count_sent, 14, sizeof(cl_mem), &b_w_count);
clSetKernelArg(k_count_sent, 15, sizeof(cl_mem), &b_total_pairs);
/* word_find_or_create: args 0-5 are constant pool buffers */
clSetKernelArg(k_word_foc, 0, sizeof(cl_mem), &b_wht_keys);
clSetKernelArg(k_word_foc, 1, sizeof(cl_mem), &b_wht_vals);
clSetKernelArg(k_word_foc, 2, sizeof(cl_mem), &b_w_hash);
clSetKernelArg(k_word_foc, 3, sizeof(cl_mem), &b_w_count);
clSetKernelArg(k_word_foc, 4, sizeof(cl_mem), &b_w_class);
clSetKernelArg(k_word_foc, 5, sizeof(cl_mem), &b_w_free);
/* compute_mi_resident: constant args */
clSetKernelArg(k_mi_res, 0, sizeof(cl_mem), &b_p_count);
clSetKernelArg(k_mi_res, 1, sizeof(cl_mem), &b_p_wa);
clSetKernelArg(k_mi_res, 2, sizeof(cl_mem), &b_p_wb);
clSetKernelArg(k_mi_res, 3, sizeof(cl_mem), &b_p_mi);
clSetKernelArg(k_mi_res, 4, sizeof(cl_mem), &b_w_count);
/* compute_mi_dirty: constant args */
clSetKernelArg(k_mi_dirty, 0, sizeof(cl_mem), &b_p_count);
clSetKernelArg(k_mi_dirty, 1, sizeof(cl_mem), &b_p_wa);
clSetKernelArg(k_mi_dirty, 2, sizeof(cl_mem), &b_p_wb);
clSetKernelArg(k_mi_dirty, 3, sizeof(cl_mem), &b_p_mi);
clSetKernelArg(k_mi_dirty, 4, sizeof(cl_mem), &b_p_flags);
clSetKernelArg(k_mi_dirty, 5, sizeof(cl_mem), &b_w_count);
/* extract_sections: constant section pool args (10-16) */
clSetKernelArg(k_extract_sec, 9, sizeof(cl_mem), &b_sht_keys);
clSetKernelArg(k_extract_sec, 10, sizeof(cl_mem), &b_sht_vals);
clSetKernelArg(k_extract_sec, 11, sizeof(cl_mem), &b_s_word);
clSetKernelArg(k_extract_sec, 12, sizeof(cl_mem), &b_s_djh);
clSetKernelArg(k_extract_sec, 13, sizeof(cl_mem), &b_s_count);
clSetKernelArg(k_extract_sec, 14, sizeof(cl_mem), &b_s_free);
clSetKernelArg(k_extract_sec, 15, sizeof(cl_mem), &b_total_secs);
/* substitute_pairs: constant args */
clSetKernelArg(k_sub_pair, 0, sizeof(cl_mem), &b_p_wa);
clSetKernelArg(k_sub_pair, 1, sizeof(cl_mem), &b_p_wb);
clSetKernelArg(k_sub_pair, 2, sizeof(cl_mem), &b_p_count);
clSetKernelArg(k_sub_pair, 3, sizeof(cl_mem), &b_p_mi);
clSetKernelArg(k_sub_pair, 4, sizeof(cl_mem), &b_p_flags);
clSetKernelArg(k_sub_pair, 5, sizeof(cl_mem), &b_w_class);
clSetKernelArg(k_sub_pair, 6, sizeof(cl_mem), &b_num_changed);
clSetKernelArg(k_sub_pair, 7, sizeof(cl_mem), &b_num_elim);
/* rebuild_pair_index: constant args */
clSetKernelArg(k_rebuild_idx, 0, sizeof(cl_mem), &b_p_wa);
clSetKernelArg(k_rebuild_idx, 1, sizeof(cl_mem), &b_p_wb);
clSetKernelArg(k_rebuild_idx, 2, sizeof(cl_mem), &b_p_count);
clSetKernelArg(k_rebuild_idx, 3, sizeof(cl_mem), &b_p_mi);
clSetKernelArg(k_rebuild_idx, 4, sizeof(cl_mem), &b_p_flags);
clSetKernelArg(k_rebuild_idx, 5, sizeof(cl_mem), &b_pht_keys);
clSetKernelArg(k_rebuild_idx, 6, sizeof(cl_mem), &b_pht_vals);
clSetKernelArg(k_rebuild_idx, 7, sizeof(cl_mem), &b_num_merged);
/* substitute_section_words: constant args */
clSetKernelArg(k_sub_sec, 0, sizeof(cl_mem), &b_s_word);
clSetKernelArg(k_sub_sec, 1, sizeof(cl_mem), &b_s_count);
clSetKernelArg(k_sub_sec, 2, sizeof(cl_mem), &b_w_class);
clSetKernelArg(k_sub_sec, 3, sizeof(cl_mem), &b_num_changed);
/* pool_stats */
clSetKernelArg(k_pool_stats, 0, sizeof(cl_mem), &b_w_free);
clSetKernelArg(k_pool_stats, 1, sizeof(cl_mem), &b_p_free);
clSetKernelArg(k_pool_stats, 2, sizeof(cl_mem), &b_s_free);
clSetKernelArg(k_pool_stats, 3, sizeof(cl_mem), &b_stats_out);
g_init = 1;
fprintf(stderr, "GPU AtomSpace ready (WORD=%dK PAIR=%dM SEC=%dM)\n",
WORD_CAP/1024, PAIR_CAP/(1024*1024), SEC_CAP/(1024*1024));
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* GPU SHUTDOWN
* ═══════════════════════════════════════════════════════════════════ */
void gpu_shutdown(void)
{
if (!g_init) return;
/* Release kernels */
cl_kernel* kerns[] = {
&k_word_foc, &k_count_sent, &k_mi_res, &k_mi_dirty,
&k_mi_stats, &k_extract_sec, &k_word_norms, &k_djh_chains,
&k_dot_prod, &k_cosines, &k_filter_cand, &k_assign_cls,
&k_sub_pair, &k_rebuild_idx, &k_sub_sec, &k_pool_stats,
&k_conn_norms, &k_conn_cosine
};
for (int i = 0; i < 18; i++)
if (*kerns[i]) clReleaseKernel(*kerns[i]);
/* Release buffers */
cl_mem* bufs[] = {
&b_wht_keys, &b_wht_vals, &b_w_hash, &b_w_count, &b_w_class, &b_w_free,
&b_pht_keys, &b_pht_vals, &b_p_wa, &b_p_wb, &b_p_count, &b_p_mi,
&b_p_flags, &b_p_free,
&b_sht_keys, &b_sht_vals, &b_s_word, &b_s_djh, &b_s_count, &b_s_free,
&b_norm_sq, &b_djh_keys, &b_djh_vals, &b_chain_next,
&b_cht_keys, &b_cht_vals, &b_c_wa, &b_c_wb, &b_c_dot, &b_c_cos, &b_c_free,
&b_total_pairs, &b_total_secs, &b_num_changed, &b_num_elim, &b_num_merged,
&b_flat_words, &b_sent_off, &b_sent_len,
&b_in_hashes, &b_out_indices,
&b_edge_p1, &b_edge_p2, &b_edge_off, &b_edge_cnt,
&b_filt_wa, &b_filt_wb, &b_filt_cos, &b_filt_cnt,
&b_stats_out,
&b_conn_matrix, &b_conn_norms,
&b_cc_wa, &b_cc_wb, &b_cc_cos, &b_cc_cnt
};
for (int i = 0; i < 55; i++)
if (*bufs[i]) clReleaseMemObject(*bufs[i]);
if (g_prog) clReleaseProgram(g_prog);
if (g_queue) clReleaseCommandQueue(g_queue);
if (g_ctx) clReleaseContext(g_ctx);
g_init = 0;
fprintf(stderr, "GPU shutdown complete\n");
}
/* ═══════════════════════════════════════════════════════════════════
* REGISTER WORDS — find-or-create in word pool
* ═══════════════════════════════════════════════════════════════════ */
int gpu_register_words(const uint64_t* name_hashes,
uint32_t* out_indices, int n)
{
if (!g_init || n <= 0) return -1;
clEnqueueWriteBuffer(g_queue, b_in_hashes, CL_FALSE, 0,
sizeof(uint64_t) * n, name_hashes, 0, NULL, NULL);
/* Set dynamic args */
clSetKernelArg(k_word_foc, 6, sizeof(cl_mem), &b_in_hashes);
clSetKernelArg(k_word_foc, 7, sizeof(cl_mem), &b_out_indices);
cl_uint un = (cl_uint)n;
clSetKernelArg(k_word_foc, 8, sizeof(cl_uint), &un);
size_t gs = round_up(n);
size_t ls = LOCAL_SIZE;
clEnqueueNDRangeKernel(g_queue, k_word_foc, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_out_indices, CL_TRUE, 0,
sizeof(uint32_t) * n, out_indices, 0, NULL, NULL);
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* COUNT SENTENCES — word pair counting within window
* ═══════════════════════════════════════════════════════════════════ */
int gpu_count_sentences(const uint32_t* flat_words,
const uint32_t* sent_offsets,
const uint32_t* sent_lengths,
int n_sentences, int total_words,
int window_size)
{
if (!g_init || total_words <= 0) return -1;
/* Upload sentence data */
clEnqueueWriteBuffer(g_queue, b_flat_words, CL_FALSE, 0,
sizeof(uint32_t) * total_words, flat_words, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_sent_off, CL_FALSE, 0,
sizeof(uint32_t) * n_sentences, sent_offsets, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_sent_len, CL_FALSE, 0,
sizeof(uint32_t) * n_sentences, sent_lengths, 0, NULL, NULL);
/* Set dynamic args */
clSetKernelArg(k_count_sent, 0, sizeof(cl_mem), &b_flat_words);
clSetKernelArg(k_count_sent, 1, sizeof(cl_mem), &b_sent_off);
clSetKernelArg(k_count_sent, 2, sizeof(cl_mem), &b_sent_len);
cl_uint ns = (cl_uint)n_sentences;
cl_uint tw = (cl_uint)total_words;
cl_uint ws = (cl_uint)window_size;
clSetKernelArg(k_count_sent, 3, sizeof(cl_uint), &ns);
clSetKernelArg(k_count_sent, 4, sizeof(cl_uint), &tw);
clSetKernelArg(k_count_sent, 5, sizeof(cl_uint), &ws);
size_t gs = round_up(total_words);
size_t ls = LOCAL_SIZE;
clEnqueueNDRangeKernel(g_queue, k_count_sent, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* COMPUTE MI — all pairs or dirty only
* ═══════════════════════════════════════════════════════════════════ */
int gpu_compute_mi(int dirty_only)
{
if (!g_init) return -1;
/* Read n_pairs from bump allocator */
uint32_t n_pairs;
clEnqueueReadBuffer(g_queue, b_p_free, CL_TRUE, 0,
sizeof(uint32_t), &n_pairs, 0, NULL, NULL);
if (n_pairs == 0) return 0;
/* Read total pair count (N for MI formula) */
uint32_t total_obs;
clEnqueueReadBuffer(g_queue, b_total_pairs, CL_TRUE, 0,
sizeof(uint32_t), &total_obs, 0, NULL, NULL);
double n_total = (double)total_obs;
if (dirty_only) {
clSetKernelArg(k_mi_dirty, 6, sizeof(double), &n_total);
cl_uint np = (cl_uint)n_pairs;
clSetKernelArg(k_mi_dirty, 7, sizeof(cl_uint), &np);
size_t gs = round_up(n_pairs);
size_t ls = LOCAL_SIZE;
clEnqueueNDRangeKernel(g_queue, k_mi_dirty, 1, NULL,
&gs, &ls, 0, NULL, NULL);
} else {
clSetKernelArg(k_mi_res, 5, sizeof(double), &n_total);
cl_uint np = (cl_uint)n_pairs;
clSetKernelArg(k_mi_res, 6, sizeof(cl_uint), &np);
size_t gs = round_up(n_pairs);
size_t ls = LOCAL_SIZE;
clEnqueueNDRangeKernel(g_queue, k_mi_res, 1, NULL,
&gs, &ls, 0, NULL, NULL);
}
clFinish(g_queue);
return (int)n_pairs;
}
/* ═══════════════════════════════════════════════════════════════════
* READ PAIR DATA — download pairs to host
* ═══════════════════════════════════════════════════════════════════ */
int gpu_read_pair_data(uint32_t* out_wa, uint32_t* out_wb,
double* out_count, double* out_mi,
int max_pairs)
{
if (!g_init) return -1;
uint32_t n_pairs;
clEnqueueReadBuffer(g_queue, b_p_free, CL_TRUE, 0,
sizeof(uint32_t), &n_pairs, 0, NULL, NULL);
int n = (int)n_pairs;
if (n > max_pairs) n = max_pairs;
if (n == 0) return 0;
clEnqueueReadBuffer(g_queue, b_p_wa, CL_FALSE, 0,
sizeof(uint32_t) * n, out_wa, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_p_wb, CL_FALSE, 0,
sizeof(uint32_t) * n, out_wb, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_p_count, CL_FALSE, 0,
sizeof(double) * n, out_count, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_p_mi, CL_TRUE, 0,
sizeof(double) * n, out_mi, 0, NULL, NULL);
return n;
}
/* ═══════════════════════════════════════════════════════════════════
* EXTRACT SECTIONS — from MST edge data
* ═══════════════════════════════════════════════════════════════════ */
int gpu_extract_sections(const uint32_t* flat_words,
const uint32_t* sent_offsets,
const uint32_t* sent_lengths,
const uint32_t* edge_p1,
const uint32_t* edge_p2,
const uint32_t* edge_offsets,
const uint32_t* edge_counts,
int n_sentences, int total_words,
int total_edges)
{
if (!g_init || total_words <= 0) return -1;
/* Upload sentence data */
clEnqueueWriteBuffer(g_queue, b_flat_words, CL_FALSE, 0,
sizeof(uint32_t) * total_words, flat_words, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_sent_off, CL_FALSE, 0,
sizeof(uint32_t) * n_sentences, sent_offsets, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_sent_len, CL_FALSE, 0,
sizeof(uint32_t) * n_sentences, sent_lengths, 0, NULL, NULL);
/* Upload edge data */
if (total_edges > 0) {
clEnqueueWriteBuffer(g_queue, b_edge_p1, CL_FALSE, 0,
sizeof(uint32_t) * total_edges, edge_p1, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_edge_p2, CL_FALSE, 0,
sizeof(uint32_t) * total_edges, edge_p2, 0, NULL, NULL);
}
clEnqueueWriteBuffer(g_queue, b_edge_off, CL_FALSE, 0,
sizeof(uint32_t) * n_sentences, edge_offsets, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_edge_cnt, CL_FALSE, 0,
sizeof(uint32_t) * n_sentences, edge_counts, 0, NULL, NULL);
/* Set args */
clSetKernelArg(k_extract_sec, 0, sizeof(cl_mem), &b_flat_words);
clSetKernelArg(k_extract_sec, 1, sizeof(cl_mem), &b_sent_off);
clSetKernelArg(k_extract_sec, 2, sizeof(cl_mem), &b_sent_len);
cl_uint ns = (cl_uint)n_sentences;
cl_uint tw = (cl_uint)total_words;
clSetKernelArg(k_extract_sec, 3, sizeof(cl_uint), &ns);
clSetKernelArg(k_extract_sec, 4, sizeof(cl_uint), &tw);
clSetKernelArg(k_extract_sec, 5, sizeof(cl_mem), &b_edge_p1);
clSetKernelArg(k_extract_sec, 6, sizeof(cl_mem), &b_edge_p2);
clSetKernelArg(k_extract_sec, 7, sizeof(cl_mem), &b_edge_off);
clSetKernelArg(k_extract_sec, 8, sizeof(cl_mem), &b_edge_cnt);
/* args 9-15 already set in init */
size_t gs = round_up(total_words);
size_t ls = LOCAL_SIZE;
clEnqueueNDRangeKernel(g_queue, k_extract_sec, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* COSINE PIPELINE — norms → chains → dots → cosines → filter
* ═══════════════════════════════════════════════════════════════════ */
int gpu_cosine_pipeline(double threshold,
uint32_t* out_wa, uint32_t* out_wb,
double* out_cos, int max_output)
{
if (!g_init) return -1;
/* Read section count */
uint32_t n_sec;
clEnqueueReadBuffer(g_queue, b_s_free, CL_TRUE, 0,
sizeof(uint32_t), &n_sec, 0, NULL, NULL);
if (n_sec == 0) return 0;
uint8_t ff = 0xFF;
uint32_t zero32 = 0;
double zero_d = 0.0;
size_t ls = LOCAL_SIZE;
/* ── Step 1: compute_word_norms ── */
clEnqueueFillBuffer(g_queue, b_norm_sq, &zero_d, sizeof(double),
0, sizeof(double) * WORD_CAP, 0, NULL, NULL);
clFinish(g_queue);
clSetKernelArg(k_word_norms, 0, sizeof(cl_mem), &b_s_word);
clSetKernelArg(k_word_norms, 1, sizeof(cl_mem), &b_s_count);
clSetKernelArg(k_word_norms, 2, sizeof(cl_mem), &b_norm_sq);
cl_uint ns = (cl_uint)n_sec;
clSetKernelArg(k_word_norms, 3, sizeof(cl_uint), &ns);
size_t gs = round_up(n_sec);
clEnqueueNDRangeKernel(g_queue, k_word_norms, 1, NULL,
&gs, &ls, 0, NULL, NULL);
/* ── Step 2: build_disjunct_chains ── */
clEnqueueFillBuffer(g_queue, b_djh_keys, &ff, 1, 0,
sizeof(uint64_t) * DJH_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_djh_vals, &ff, 1, 0,
sizeof(uint32_t) * DJH_HT_CAP, 0, NULL, NULL);
clFinish(g_queue);
clSetKernelArg(k_djh_chains, 0, sizeof(cl_mem), &b_s_djh);
clSetKernelArg(k_djh_chains, 1, sizeof(cl_mem), &b_s_count);
clSetKernelArg(k_djh_chains, 2, sizeof(cl_mem), &b_djh_keys);
clSetKernelArg(k_djh_chains, 3, sizeof(cl_mem), &b_djh_vals);
clSetKernelArg(k_djh_chains, 4, sizeof(cl_mem), &b_chain_next);
clSetKernelArg(k_djh_chains, 5, sizeof(cl_uint), &ns);
clEnqueueNDRangeKernel(g_queue, k_djh_chains, 1, NULL,
&gs, &ls, 0, NULL, NULL);
/* ── Step 3: accumulate_dot_products ── */
clEnqueueFillBuffer(g_queue, b_cht_keys, &ff, 1, 0,
sizeof(uint64_t) * CAND_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_cht_vals, &ff, 1, 0,
sizeof(uint32_t) * CAND_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_c_free, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clFinish(g_queue);
clSetKernelArg(k_dot_prod, 0, sizeof(cl_mem), &b_s_word);
clSetKernelArg(k_dot_prod, 1, sizeof(cl_mem), &b_s_djh);
clSetKernelArg(k_dot_prod, 2, sizeof(cl_mem), &b_s_count);
clSetKernelArg(k_dot_prod, 3, sizeof(cl_mem), &b_djh_keys);
clSetKernelArg(k_dot_prod, 4, sizeof(cl_mem), &b_djh_vals);
clSetKernelArg(k_dot_prod, 5, sizeof(cl_mem), &b_chain_next);
clSetKernelArg(k_dot_prod, 6, sizeof(cl_mem), &b_cht_keys);
clSetKernelArg(k_dot_prod, 7, sizeof(cl_mem), &b_cht_vals);
clSetKernelArg(k_dot_prod, 8, sizeof(cl_mem), &b_c_wa);
clSetKernelArg(k_dot_prod, 9, sizeof(cl_mem), &b_c_wb);
clSetKernelArg(k_dot_prod, 10, sizeof(cl_mem), &b_c_dot);
clSetKernelArg(k_dot_prod, 11, sizeof(cl_mem), &b_c_free);
clSetKernelArg(k_dot_prod, 12, sizeof(cl_uint), &ns);
clEnqueueNDRangeKernel(g_queue, k_dot_prod, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
/* Read candidate count */
uint32_t n_cand;
clEnqueueReadBuffer(g_queue, b_c_free, CL_TRUE, 0,
sizeof(uint32_t), &n_cand, 0, NULL, NULL);
if (n_cand == 0) return 0;
/* ── Step 4: compute_cosines ── */
clSetKernelArg(k_cosines, 0, sizeof(cl_mem), &b_c_wa);
clSetKernelArg(k_cosines, 1, sizeof(cl_mem), &b_c_wb);
clSetKernelArg(k_cosines, 2, sizeof(cl_mem), &b_c_dot);
clSetKernelArg(k_cosines, 3, sizeof(cl_mem), &b_c_cos);
clSetKernelArg(k_cosines, 4, sizeof(cl_mem), &b_norm_sq);
cl_uint nc = (cl_uint)n_cand;
clSetKernelArg(k_cosines, 5, sizeof(cl_uint), &nc);
gs = round_up(n_cand);
clEnqueueNDRangeKernel(g_queue, k_cosines, 1, NULL,
&gs, &ls, 0, NULL, NULL);
/* ── Step 5: filter_candidates ── */
clEnqueueFillBuffer(g_queue, b_filt_cnt, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clFinish(g_queue);
clSetKernelArg(k_filter_cand, 0, sizeof(cl_mem), &b_c_wa);
clSetKernelArg(k_filter_cand, 1, sizeof(cl_mem), &b_c_wb);
clSetKernelArg(k_filter_cand, 2, sizeof(cl_mem), &b_c_cos);
clSetKernelArg(k_filter_cand, 3, sizeof(cl_uint), &nc);
clSetKernelArg(k_filter_cand, 4, sizeof(double), &threshold);
clSetKernelArg(k_filter_cand, 5, sizeof(cl_mem), &b_filt_wa);
clSetKernelArg(k_filter_cand, 6, sizeof(cl_mem), &b_filt_wb);
clSetKernelArg(k_filter_cand, 7, sizeof(cl_mem), &b_filt_cos);
clSetKernelArg(k_filter_cand, 8, sizeof(cl_mem), &b_filt_cnt);
cl_uint mo = (cl_uint)(max_output < MAX_FILTER_OUT ?
max_output : MAX_FILTER_OUT);
clSetKernelArg(k_filter_cand, 9, sizeof(cl_uint), &mo);
clEnqueueNDRangeKernel(g_queue, k_filter_cand, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
/* Read results */
uint32_t n_filt;
clEnqueueReadBuffer(g_queue, b_filt_cnt, CL_TRUE, 0,
sizeof(uint32_t), &n_filt, 0, NULL, NULL);
int n_out = (int)n_filt;
if (n_out > max_output) n_out = max_output;
if (n_out > (int)MAX_FILTER_OUT) n_out = MAX_FILTER_OUT;
if (n_out == 0) return 0;
clEnqueueReadBuffer(g_queue, b_filt_wa, CL_FALSE, 0,
sizeof(uint32_t) * n_out, out_wa, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_filt_wb, CL_FALSE, 0,
sizeof(uint32_t) * n_out, out_wb, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_filt_cos, CL_TRUE, 0,
sizeof(double) * n_out, out_cos, 0, NULL, NULL);
return n_out;
}
/* ═══════════════════════════════════════════════════════════════════
* BULK LOAD SECTIONS — import from CogServer without re-counting
*
* Writes directly into section pool arrays (b_s_word, b_s_djh,
* b_s_count). Replaces any existing sections. Used to unify the
* GPU data store with the CogServer's existing section data.
*
* Args:
* word_ids — GPU pool index per section (from gpu_register_words)
* djh_hashes — FNV1a hash of disjunct string per section
* counts — observation count per section
* n — number of sections to load
*
* Returns 0 on success, -1 on error.
* ═══════════════════════════════════════════════════════════════════ */
int gpu_bulk_load_sections(const uint32_t* word_ids,
const uint64_t* djh_hashes,
const double* counts, int n)
{
if (!g_init || n <= 0) return -1;
if (n > SEC_CAP) {
fprintf(stderr, "gpu_bulk_load_sections: %d > SEC_CAP (%d)\n",
n, SEC_CAP);
n = SEC_CAP;
}
/* Upload section data directly to GPU buffers */
clEnqueueWriteBuffer(g_queue, b_s_word, CL_FALSE, 0,
sizeof(uint32_t) * n, word_ids, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_s_djh, CL_FALSE, 0,
sizeof(uint64_t) * n, djh_hashes, 0, NULL, NULL);
clEnqueueWriteBuffer(g_queue, b_s_count, CL_FALSE, 0,
sizeof(double) * n, counts, 0, NULL, NULL);
/* Set section count (bump allocator) */
uint32_t un = (uint32_t)n;
clEnqueueWriteBuffer(g_queue, b_s_free, CL_TRUE, 0,
sizeof(uint32_t), &un, 0, NULL, NULL);
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* ASSIGN CLASSES
* ═══════════════════════════════════════════════════════════════════ */
int gpu_assign_classes(const uint32_t* word_ids,
const uint32_t* class_ids, int n)
{
if (!g_init || n <= 0) return -1;
cl_mem d_wids = clCreateBuffer(g_ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(uint32_t) * n, (void*)word_ids, NULL);
cl_mem d_cids = clCreateBuffer(g_ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(uint32_t) * n, (void*)class_ids, NULL);
clSetKernelArg(k_assign_cls, 0, sizeof(cl_mem), &b_w_class);
clSetKernelArg(k_assign_cls, 1, sizeof(cl_mem), &d_wids);
clSetKernelArg(k_assign_cls, 2, sizeof(cl_mem), &d_cids);
cl_uint un = (cl_uint)n;
clSetKernelArg(k_assign_cls, 3, sizeof(cl_uint), &un);
size_t gs = round_up(n);
size_t ls = LOCAL_SIZE;
clEnqueueNDRangeKernel(g_queue, k_assign_cls, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
clReleaseMemObject(d_wids);
clReleaseMemObject(d_cids);
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* SUBSTITUTE — pairs + rebuild HT + sections
* ═══════════════════════════════════════════════════════════════════ */
int gpu_substitute(uint32_t* out_changed, uint32_t* out_eliminated,
uint32_t* out_merged)
{
if (!g_init) return -1;
/* Read pool sizes */
uint32_t n_pairs, n_sec;
clEnqueueReadBuffer(g_queue, b_p_free, CL_TRUE, 0,
sizeof(uint32_t), &n_pairs, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_s_free, CL_TRUE, 0,
sizeof(uint32_t), &n_sec, 0, NULL, NULL);
uint32_t zero32 = 0;
uint8_t ff = 0xFF;
size_t ls = LOCAL_SIZE;
/* Reset counters */
clEnqueueFillBuffer(g_queue, b_num_changed, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_num_elim, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_num_merged, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clFinish(g_queue);
/* Step 1: substitute_pairs */
if (n_pairs > 0) {
cl_uint np = (cl_uint)n_pairs;
clSetKernelArg(k_sub_pair, 8, sizeof(cl_uint), &np);
size_t gs = round_up(n_pairs);
clEnqueueNDRangeKernel(g_queue, k_sub_pair, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
/* Step 2: clear and rebuild pair HT */
clEnqueueFillBuffer(g_queue, b_pht_keys, &ff, 1, 0,
sizeof(uint64_t) * PAIR_HT_CAP, 0, NULL, NULL);
clEnqueueFillBuffer(g_queue, b_pht_vals, &ff, 1, 0,
sizeof(uint32_t) * PAIR_HT_CAP, 0, NULL, NULL);
clFinish(g_queue);
clSetKernelArg(k_rebuild_idx, 8, sizeof(cl_uint), &np);
clEnqueueNDRangeKernel(g_queue, k_rebuild_idx, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
}
/* Step 3: substitute_section_words */
if (n_sec > 0) {
/* Reset num_changed for section count */
clEnqueueFillBuffer(g_queue, b_num_changed, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clFinish(g_queue);
cl_uint ns = (cl_uint)n_sec;
clSetKernelArg(k_sub_sec, 4, sizeof(cl_uint), &ns);
size_t gs = round_up(n_sec);
clEnqueueNDRangeKernel(g_queue, k_sub_sec, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
}
/* Read results */
if (out_changed)
clEnqueueReadBuffer(g_queue, b_num_changed, CL_TRUE, 0,
sizeof(uint32_t), out_changed, 0, NULL, NULL);
if (out_eliminated)
clEnqueueReadBuffer(g_queue, b_num_elim, CL_TRUE, 0,
sizeof(uint32_t), out_eliminated, 0, NULL, NULL);
if (out_merged)
clEnqueueReadBuffer(g_queue, b_num_merged, CL_TRUE, 0,
sizeof(uint32_t), out_merged, 0, NULL, NULL);
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* STATISTICS
* ═══════════════════════════════════════════════════════════════════ */
int gpu_get_stats(uint32_t* n_words, uint32_t* n_pairs,
uint32_t* n_sections)
{
if (!g_init) return -1;
/* Run pool_stats kernel (single thread) */
size_t gs = 1, ls = 1;
clEnqueueNDRangeKernel(g_queue, k_pool_stats, 1, NULL,
&gs, &ls, 0, NULL, NULL);
uint32_t stats[4];
clEnqueueReadBuffer(g_queue, b_stats_out, CL_TRUE, 0,
sizeof(uint32_t) * 3, stats, 0, NULL, NULL);
if (n_words) *n_words = stats[0];
if (n_pairs) *n_pairs = stats[1];
if (n_sections) *n_sections = stats[2];
return 0;
}
/* ═══════════════════════════════════════════════════════════════════
* GET TOTAL PAIR OBSERVATIONS (N for MI formula)
* ═══════════════════════════════════════════════════════════════════ */
uint32_t gpu_get_total_observations(void)
{
if (!g_init) return 0;
uint32_t total;
clEnqueueReadBuffer(g_queue, b_total_pairs, CL_TRUE, 0,
sizeof(uint32_t), &total, 0, NULL, NULL);
return total;
}
/* ═══════════════════════════════════════════════════════════════════
* CONNECTOR COSINE — brute-force pairwise on dense matrix
*
* Input: dense float matrix [n_words × n_conn] in row-major order.
* Each row = one word's connector count vector.
* Output: pairs (wa, wb, cosine) above threshold, sorted by cosine desc.
*
* Returns number of output pairs, or -1 on error.
* ═══════════════════════════════════════════════════════════════════ */
int gpu_connector_cosine(const float* matrix, int n_words, int n_conn,
float threshold,
uint32_t* out_wa, uint32_t* out_wb,
float* out_cos, int max_output)
{
if (!g_init) return -1;
if (n_words <= 1 || n_conn <= 0) return 0;
if (n_words > CONN_WORD_CAP) {
fprintf(stderr, "gpu_connector_cosine: n_words %d > cap %d\n",
n_words, CONN_WORD_CAP);
n_words = CONN_WORD_CAP;
}
if (n_conn > CONN_DIM_CAP) {
fprintf(stderr, "gpu_connector_cosine: n_conn %d > cap %d\n",
n_conn, CONN_DIM_CAP);
n_conn = CONN_DIM_CAP;
}
int mo = max_output < CONN_MAX_OUTPUT ? max_output : CONN_MAX_OUTPUT;
cl_int err;
size_t ls = LOCAL_SIZE;
uint32_t zero32 = 0;
/* Upload dense matrix */
size_t matrix_bytes = sizeof(float) * n_words * n_conn;
err = clEnqueueWriteBuffer(g_queue, b_conn_matrix, CL_FALSE, 0,
matrix_bytes, matrix, 0, NULL, NULL);
if (err != CL_SUCCESS) {
fprintf(stderr, "gpu_connector_cosine: upload failed %d\n", err);
return -1;
}
/* Step 1: compute norms */
cl_uint uw = (cl_uint)n_words;
cl_uint uc = (cl_uint)n_conn;
clSetKernelArg(k_conn_norms, 0, sizeof(cl_mem), &b_conn_matrix);
clSetKernelArg(k_conn_norms, 1, sizeof(cl_mem), &b_conn_norms);
clSetKernelArg(k_conn_norms, 2, sizeof(cl_uint), &uw);
clSetKernelArg(k_conn_norms, 3, sizeof(cl_uint), &uc);
size_t gs = round_up(n_words);
clEnqueueNDRangeKernel(g_queue, k_conn_norms, 1, NULL,
&gs, &ls, 0, NULL, NULL);
/* Step 2: brute-force pairwise cosine */
clEnqueueFillBuffer(g_queue, b_cc_cnt, &zero32, sizeof(uint32_t),
0, sizeof(uint32_t), 0, NULL, NULL);
clFinish(g_queue);
cl_uint umo = (cl_uint)mo;
clSetKernelArg(k_conn_cosine, 0, sizeof(cl_mem), &b_conn_matrix);
clSetKernelArg(k_conn_cosine, 1, sizeof(cl_mem), &b_conn_norms);
clSetKernelArg(k_conn_cosine, 2, sizeof(cl_uint), &uw);
clSetKernelArg(k_conn_cosine, 3, sizeof(cl_uint), &uc);
clSetKernelArg(k_conn_cosine, 4, sizeof(float), &threshold);
clSetKernelArg(k_conn_cosine, 5, sizeof(cl_mem), &b_cc_wa);
clSetKernelArg(k_conn_cosine, 6, sizeof(cl_mem), &b_cc_wb);
clSetKernelArg(k_conn_cosine, 7, sizeof(cl_mem), &b_cc_cos);
clSetKernelArg(k_conn_cosine, 8, sizeof(cl_mem), &b_cc_cnt);
clSetKernelArg(k_conn_cosine, 9, sizeof(cl_uint), &umo);
uint32_t total_pairs = (n_words * (n_words - 1)) / 2;
gs = round_up(total_pairs);
clEnqueueNDRangeKernel(g_queue, k_conn_cosine, 1, NULL,
&gs, &ls, 0, NULL, NULL);
clFinish(g_queue);
/* Read results */
uint32_t n_out;
clEnqueueReadBuffer(g_queue, b_cc_cnt, CL_TRUE, 0,
sizeof(uint32_t), &n_out, 0, NULL, NULL);
int n = (int)n_out;
if (n > mo) n = mo;
if (n == 0) return 0;
clEnqueueReadBuffer(g_queue, b_cc_wa, CL_FALSE, 0,
sizeof(uint32_t) * n, out_wa, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_cc_wb, CL_FALSE, 0,
sizeof(uint32_t) * n, out_wb, 0, NULL, NULL);
clEnqueueReadBuffer(g_queue, b_cc_cos, CL_TRUE, 0,
sizeof(float) * n, out_cos, 0, NULL, NULL);
return n;
}
/* ═══════════════════════════════════════════════════════════════════
* TEST MAIN
* ═══════════════════════════════════════════════════════════════════ */
#ifdef TEST_MAIN
int main(int argc, char** argv)
{
const char* kdir = argc > 1 ? argv[1] :
"/storage/projects/code/hypercore-kb/opencog-kb/repos/"
"atomese-simd/opencog/opencl/atomspace";
printf("=== GPU AtomSpace Bridge Test ===\n\n");
/* Init */
if (gpu_init(kdir) != 0) {
printf("FAIL: gpu_init\n");
return 1;
}
/* Register words */
uint64_t hashes[] = {100, 200, 300, 400, 500};
uint32_t indices[5];
gpu_register_words(hashes, indices, 5);
printf("Words registered: %u %u %u %u %u\n",
indices[0], indices[1], indices[2], indices[3], indices[4]);
/* Count sentences */
uint32_t flat[] = {indices[0], indices[1], indices[2],
indices[3], indices[4]};
uint32_t off[] = {0, 3};
uint32_t len[] = {3, 2};
gpu_count_sentences(flat, off, len, 2, 5, 6);
/* Stats */
uint32_t nw, np, ns;
gpu_get_stats(&nw, &np, &ns);
printf("Words: %u Pairs: %u Sections: %u\n", nw, np, ns);
/* MI */
int mi_pairs = gpu_compute_mi(0);
printf("MI computed for %d pairs\n", mi_pairs);
/* Read pairs */
uint32_t wa[64], wb[64];
double cnt[64], mi[64];
int n = gpu_read_pair_data(wa, wb, cnt, mi, 64);
printf("Read %d pairs:\n", n);
for (int i = 0; i < n && i < 10; i++) {
printf(" (%u, %u) count=%.1f mi=%.4f\n",
wa[i], wb[i], cnt[i], mi[i]);
}
/* Assign classes: words 0,1 → class 100 */
uint32_t cwids[] = {indices[0], indices[1]};
uint32_t ccids[] = {100, 100};
gpu_assign_classes(cwids, ccids, 2);
/* Substitute */
uint32_t changed, elim, merged;
gpu_substitute(&changed, &elim, &merged);
printf("Substitute: changed=%u eliminated=%u merged=%u\n",
changed, elim, merged);
/* Final stats */
gpu_get_stats(&nw, &np, &ns);
printf("Final: Words=%u Pairs=%u Sections=%u\n", nw, np, ns);
gpu_shutdown();
printf("\n=== PASS ===\n");
return 0;
}
#endif