/**
* tmfg-kernel.cl - Triangulated Maximally Filtered Graph on OpenCL
*
* TMFG builds a maximally planar graph by inserting nodes into triangles.
* This avoids planarity testing entirely - the result is planar by construction.
*
* For linguistic analysis:
* - Nodes = words
* - Edge weights = Mutual Information (MI)
* - Output = grammatical skeleton (no crossing dependencies)
*/
/**
* find_initial_triangle - Find the 3 words with highest pairwise MI
*
* This forms the seed triangle for TMFG construction.
* Run with 1 work-item (reduction operation).
*/
__kernel void find_initial_triangle(
__global const double* mi_matrix,
__global int* best_triangle, // Output: [v1, v2, v3]
__global double* best_score, // Output: total MI
const int vocab_size,
const unsigned long sz) // = vocab_size
{
// Single thread does exhaustive search for best triangle
// (Small vocab makes this fast; could parallelize if needed)
if (get_global_id(0) != 0) return;
double best = -1e30;
int best_v1 = 0, best_v2 = 1, best_v3 = 2;
for (int v1 = 0; v1 < vocab_size - 2; v1++) {
for (int v2 = v1 + 1; v2 < vocab_size - 1; v2++) {
double mi_12 = mi_matrix[v1 * vocab_size + v2];
for (int v3 = v2 + 1; v3 < vocab_size; v3++) {
double mi_13 = mi_matrix[v1 * vocab_size + v3];
double mi_23 = mi_matrix[v2 * vocab_size + v3];
double score = mi_12 + mi_13 + mi_23;
if (score > best) {
best = score;
best_v1 = v1;
best_v2 = v2;
best_v3 = v3;
}
}
}
}
best_triangle[0] = best_v1;
best_triangle[1] = best_v2;
best_triangle[2] = best_v3;
*best_score = best;
}
/**
* score_node_triangles - For each uninserted node, find best triangle to insert into
*
* Each work-item handles one node, scanning all triangles.
* This is the parallelizable core of TMFG.
*/
__kernel void score_node_triangles(
__global const double* mi_matrix,
__global const int* triangles, // Current triangles [t0_v1, t0_v2, t0_v3, t1_v1, ...]
__global const int* node_inserted, // 1 if node already in graph, 0 otherwise
__global int* best_triangle_per_node,
__global double* best_score_per_node,
const int n_triangles,
const int vocab_size,
const unsigned long sz)
{
size_t node = get_global_id(0);
if (node >= sz) return;
// Skip already-inserted nodes
if (node_inserted[node]) {
best_triangle_per_node[node] = -1;
best_score_per_node[node] = -1e30;
return;
}
double best = -1e30;
int best_tri = -1;
// Find triangle with highest total MI to this node
for (int t = 0; t < n_triangles; t++) {
int v1 = triangles[t * 3];
int v2 = triangles[t * 3 + 1];
int v3 = triangles[t * 3 + 2];
// Skip degenerate triangles
if (v1 < 0 || v2 < 0 || v3 < 0) continue;
double mi_1 = mi_matrix[node * vocab_size + v1];
double mi_2 = mi_matrix[node * vocab_size + v2];
double mi_3 = mi_matrix[node * vocab_size + v3];
double score = mi_1 + mi_2 + mi_3;
if (score > best) {
best = score;
best_tri = t;
}
}
best_triangle_per_node[node] = best_tri;
best_score_per_node[node] = best;
}
/**
* find_best_node_to_insert - Reduction to find globally best node
*
* After score_node_triangles, find which node has the highest score.
*/
__kernel void find_best_node_to_insert(
__global const double* best_score_per_node,
__global const int* node_inserted,
__global int* best_node,
__global double* best_score,
const unsigned long sz)
{
// Parallel reduction
__local double local_scores[256];
__local int local_nodes[256];
size_t lid = get_local_id(0);
size_t gid = get_global_id(0);
// Load into local memory
if (gid < sz && !node_inserted[gid]) {
local_scores[lid] = best_score_per_node[gid];
local_nodes[lid] = gid;
} else {
local_scores[lid] = -1e30;
local_nodes[lid] = -1;
}
barrier(CLK_LOCAL_MEM_FENCE);
// Reduce within work-group
for (int stride = get_local_size(0) / 2; stride > 0; stride /= 2) {
if (lid < stride) {
if (local_scores[lid + stride] > local_scores[lid]) {
local_scores[lid] = local_scores[lid + stride];
local_nodes[lid] = local_nodes[lid + stride];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
}
// First thread writes result
if (lid == 0) {
// Atomic max (simplified - real impl needs atomic compare-exchange)
if (local_scores[0] > *best_score) {
*best_score = local_scores[0];
*best_node = local_nodes[0];
}
}
}
/**
* insert_node_into_triangle - Update triangle list after insertion
*
* When node N is inserted into triangle (A,B,C), we:
* 1. Remove triangle (A,B,C)
* 2. Add triangles (N,A,B), (N,B,C), (N,A,C)
*
* Run with 1 work-item per insertion (sequential step).
*/
__kernel void insert_node_into_triangle(
__global int* triangles,
__global int* n_triangles,
__global int* node_inserted,
__global int* edges, // Output: edges added [src, dst, ...]
__global int* n_edges,
const int node_to_insert,
const int triangle_idx,
const unsigned long sz)
{
if (get_global_id(0) != 0) return;
int v1 = triangles[triangle_idx * 3];
int v2 = triangles[triangle_idx * 3 + 1];
int v3 = triangles[triangle_idx * 3 + 2];
// Mark node as inserted
node_inserted[node_to_insert] = 1;
// Replace old triangle with first new triangle
triangles[triangle_idx * 3] = node_to_insert;
triangles[triangle_idx * 3 + 1] = v1;
triangles[triangle_idx * 3 + 2] = v2;
// Add two more triangles
int nt = *n_triangles;
triangles[nt * 3] = node_to_insert;
triangles[nt * 3 + 1] = v2;
triangles[nt * 3 + 2] = v3;
triangles[(nt + 1) * 3] = node_to_insert;
triangles[(nt + 1) * 3 + 1] = v1;
triangles[(nt + 1) * 3 + 2] = v3;
*n_triangles = nt + 2;
// Record edges added
int ne = *n_edges;
edges[ne * 2] = node_to_insert;
edges[ne * 2 + 1] = v1;
edges[(ne + 1) * 2] = node_to_insert;
edges[(ne + 1) * 2 + 1] = v2;
edges[(ne + 2) * 2] = node_to_insert;
edges[(ne + 2) * 2 + 1] = v3;
*n_edges = ne + 3;
}
/**
* batch_sentence_tmfg - Process multiple sentences in parallel
*
* For linguistic analysis, we run TMFG on each sentence independently.
* Each work-item handles one sentence.
*/
__kernel void batch_sentence_tmfg(
__global const double* mi_matrix, // Global MI matrix [vocab × vocab]
__global const int* sentences, // Word indices [max_len × n_sentences]
__global const int* sentence_lens,
__global int* tmfg_edges, // Output: edges per sentence
__global int* tmfg_edge_counts,
const int vocab_size,
const int max_len,
const unsigned long sz)
{
size_t sent_idx = get_global_id(0);
if (sent_idx >= sz) return;
int n = sentence_lens[sent_idx];
if (n < 3) {
tmfg_edge_counts[sent_idx] = 0;
return;
}
__global const int* words = sentences + sent_idx * max_len;
__global int* edges = tmfg_edges + sent_idx * (3 * max_len - 6) * 2;
// Local arrays for this sentence's TMFG
int inserted[32]; // Max 32 words per sentence
int triangles[90]; // Max 30 triangles × 3 vertices
int n_tri = 0;
int n_edges = 0;
for (int i = 0; i < n; i++) inserted[i] = 0;
// Step 1: Find initial triangle (best 3 words)
double best_tri_score = -1e30;
int best_v1 = 0, best_v2 = 1, best_v3 = 2;
for (int i = 0; i < n - 2; i++) {
int w1 = words[i];
for (int j = i + 1; j < n - 1; j++) {
int w2 = words[j];
double mi_12 = mi_matrix[w1 * vocab_size + w2];
for (int k = j + 1; k < n; k++) {
int w3 = words[k];
double score = mi_12
+ mi_matrix[w1 * vocab_size + w3]
+ mi_matrix[w2 * vocab_size + w3];
if (score > best_tri_score) {
best_tri_score = score;
best_v1 = i; best_v2 = j; best_v3 = k;
}
}
}
}
// Initialize with first triangle
triangles[0] = best_v1;
triangles[1] = best_v2;
triangles[2] = best_v3;
n_tri = 1;
inserted[best_v1] = 1;
inserted[best_v2] = 1;
inserted[best_v3] = 1;
// Record initial edges
edges[0] = best_v1; edges[1] = best_v2;
edges[2] = best_v1; edges[3] = best_v3;
edges[4] = best_v2; edges[5] = best_v3;
n_edges = 3;
// Step 2: Insert remaining nodes
for (int iter = 0; iter < n - 3; iter++) {
// Find best node to insert
double best_score = -1e30;
int best_node = -1;
int best_tri_idx = -1;
for (int node = 0; node < n; node++) {
if (inserted[node]) continue;
int w_node = words[node];
for (int t = 0; t < n_tri; t++) {
int v1 = triangles[t * 3];
int v2 = triangles[t * 3 + 1];
int v3 = triangles[t * 3 + 2];
double score = mi_matrix[w_node * vocab_size + words[v1]]
+ mi_matrix[w_node * vocab_size + words[v2]]
+ mi_matrix[w_node * vocab_size + words[v3]];
if (score > best_score) {
best_score = score;
best_node = node;
best_tri_idx = t;
}
}
}
if (best_node < 0) break;
// Insert node into triangle
int v1 = triangles[best_tri_idx * 3];
int v2 = triangles[best_tri_idx * 3 + 1];
int v3 = triangles[best_tri_idx * 3 + 2];
inserted[best_node] = 1;
// Replace old triangle and add two new ones
triangles[best_tri_idx * 3] = best_node;
triangles[best_tri_idx * 3 + 1] = v1;
triangles[best_tri_idx * 3 + 2] = v2;
triangles[n_tri * 3] = best_node;
triangles[n_tri * 3 + 1] = v2;
triangles[n_tri * 3 + 2] = v3;
triangles[(n_tri + 1) * 3] = best_node;
triangles[(n_tri + 1) * 3 + 1] = v1;
triangles[(n_tri + 1) * 3 + 2] = v3;
n_tri += 2;
// Record new edges
edges[n_edges * 2] = best_node;
edges[n_edges * 2 + 1] = v1;
edges[(n_edges + 1) * 2] = best_node;
edges[(n_edges + 1) * 2 + 1] = v2;
edges[(n_edges + 2) * 2] = best_node;
edges[(n_edges + 2) * 2 + 1] = v3;
n_edges += 3;
}
tmfg_edge_counts[sent_idx] = n_edges;
}
/**
* extract_sections - Convert TMFG edges to Sections (word + connector-set)
*
* For each word, count left and right connections based on position.
*/
__kernel void extract_sections(
__global const int* sentences,
__global const int* sentence_lens,
__global const int* tmfg_edges,
__global const int* tmfg_edge_counts,
__global int* left_counts, // Output: left connector count per word
__global int* right_counts, // Output: right connector count per word
const int max_len,
const unsigned long sz)
{
size_t sent_idx = get_global_id(0);
if (sent_idx >= sz) return;
int n = sentence_lens[sent_idx];
int n_edges = tmfg_edge_counts[sent_idx];
__global const int* edges = tmfg_edges + sent_idx * (3 * max_len - 6) * 2;
__global int* left = left_counts + sent_idx * max_len;
__global int* right = right_counts + sent_idx * max_len;
// Initialize counts
for (int i = 0; i < n; i++) {
left[i] = 0;
right[i] = 0;
}
// Process edges
for (int e = 0; e < n_edges; e++) {
int pos1 = edges[e * 2];
int pos2 = edges[e * 2 + 1];
if (pos1 < pos2) {
right[pos1]++; // pos1 connects right to pos2
left[pos2]++; // pos2 connects left to pos1
} else {
right[pos2]++;
left[pos1]++;
}
}
}