Task: bellman_ford
Category: Graph algorithms
Difficulty: 2
Tags: request_based,graph,sssp,CSR_graph
Source: Boost Graph Library: https://github.com/boostorg/graph
Description: Request-based Bellman--Ford single-source shortest-path computation on CSR graphs.

Prompt template (YAML):
# Bellman-Ford SSSP — Prompt Template
#
# Used by framework/generate_prompt.py to assemble level-specific prompts.
# The framework auto-injects: generic constraints, input sizes (from task.json),
# and cpu_reference.c content.

task_description: |
  Given a weighted directed graph (CSR format) and a batch of `(source, target)` queries,
  compute the shortest distance from each source to its target.

interface: |
  ```c
  // Single end-to-end call. Do H2D + kernel + D2H and synchronize before returning.
  // Called multiple times (warmup + timed); must be idempotent.
  extern "C" void solution_compute(
      int V,
      int E,
      const int* h_row_offsets,
      const int* h_col_indices,
      const float* h_weights,
      int num_requests,
      const int* h_sources,
      const int* h_targets,
      float* h_distances
  );
  ```

input_size_notes: |
  Queries are 100 `(source, target)` pairs: 10 distinct sources × 10 targets each.

output_contract: |
  - `h_distances[i]` = shortest distance from `h_sources[i]` to `h_targets[i]`
  - Unreachable: set to `1e30f` (NOT `INFINITY`, NOT `-1`)
  - Tolerance: results within `atol=0.1` of CPU reference are accepted
  - `solution_compute` must **synchronize** before returning
    (all results must be visible in `h_distances` when the function returns)
  - Queries with the same source are grouped consecutively (10 targets per source)

# ── L1 only: algorithm background ──

algorithm_background: |
  **Bellman-Ford** computes single-source shortest paths in a weighted directed graph
  stored in CSR (Compressed Sparse Row) format:
  - `row_offsets[V+1]`: node u's edges are at indices `row_offsets[u]` to `row_offsets[u+1]-1`
  - `col_indices[E]`: destination node for each edge
  - `weights[E]`: edge weight (positive floats)

  The algorithm runs up to `V-1` rounds. Each round relaxes all edges:
  if `dist[u] + w(u,v) < dist[v]`, update `dist[v]`. Stop early if no update occurs in a round.

# ── Hints by level ──

hints_l2: |
  - Copy graph data to GPU **once** inside `solution_compute`; keep state local to each call
  - Queries with the same `source` can share a single Bellman-Ford run (deduplicate sources)
  - Use early termination: stop iterating when no edge is relaxed in a round
  - Consider edge-parallel or vertex-parallel relaxation strategies

hints_l1: |
  ### 1. Persistent GPU Graph Data
  Copy `row_offsets`, `col_indices`, and `weights` to device memory inside `solution_compute`.
  Also pre-allocate a `dist` buffer on the device. keep state local to each call — no per-call allocation.

  ### 2. Source Deduplication
  The 100 queries have only 10 distinct sources. Group queries by source and run
  Bellman-Ford once per unique source. After computing the full `dist[]` array for
  a source, look up `dist[target]` for all 10 targets of that source.

  ### 3. Edge-Parallel Relaxation
  Instead of node-parallel (one thread per node iterating over its neighbors),
  use edge-parallel: assign one thread per edge `(u,v)`. Each thread reads `dist[u]`,
  computes `dist[u] + weight`, and atomically updates `dist[v]` using an
  `atomicMin`-equivalent CAS loop for floats:

  ```cuda
  __device__ void atomicMinFloat(float* addr, float val) {
      int* addr_as_int = (int*)addr;
      int old = *addr_as_int, assumed;
      do {
          assumed = old;
          if (__int_as_float(assumed) <= val) break;
          old = atomicCAS(addr_as_int, assumed, __float_as_int(val));
      } while (assumed != old);
  }
  ```

  ### 4. Early Termination
  Use a device-side `int* d_updated` flag. Clear it with `cudaMemset` before each
  round. Any thread that performs a successful relaxation sets `*d_updated = 1`.
  Copy back to host after each round; if zero, break out of the iteration loop.

  ### 5. Convergence Optimization
  For sparse graphs, most nodes converge in far fewer than `V-1` rounds.
  The early termination check makes the GPU version much faster in practice,
  especially on the `large` input.

  ### 6. Memory Coalescing
  When using edge-parallel, ensure threads in a warp access consecutive edges
  (consecutive entries in `col_indices` and `weights`), which naturally follows
  the CSR layout.



CPU reference implementation (C):
// cpu_reference.c — bellman_ford CPU baseline
//
// Implements the two LLM-facing functions:
//   solution_init / solution_compute
//
// NO file I/O here. All I/O is handled by task_io_cpu.c.
//
// Build (via task_io + harness):
//   gcc -O2 -I framework/
//       framework/harness_cpu.c tasks/bellman_ford/task_io_cpu.c
//       tasks/bellman_ford/cpu_reference.c -o solution_cpu -lm

#include <stdlib.h>
#include <string.h>

#define INF_VAL 1e30f

// ===== Module-level state =====
static int g_V, g_E;
static const int*   g_row_offsets;
static const int*   g_col_indices;
static const float* g_weights;
static float*       g_dist_buf;   // working buffer, size V

static void _orbench_old_init(int V, int E,
                   const int* h_row_offsets,
                   const int* h_col_indices,
                   const float* h_weights) {
    g_V = V;
    g_E = E;
    g_row_offsets = h_row_offsets;
    g_col_indices = h_col_indices;
    g_weights     = h_weights;
    g_dist_buf    = (float*)malloc((size_t)V * sizeof(float));
}

static void _orbench_old_compute(int num_requests,
                      const int* sources,
                      const int* targets,
                      float* distances) {
    for (int r = 0; r < num_requests; r++) {
        int src = sources[r];
        int tgt = targets[r];

        // Bellman-Ford from src
        for (int i = 0; i < g_V; i++) g_dist_buf[i] = INF_VAL;
        if (src >= 0 && src < g_V) g_dist_buf[src] = 0.0f;

        for (int round = 0; round < g_V - 1; round++) {
            int updated = 0;
            for (int u = 0; u < g_V; u++) {
                float du = g_dist_buf[u];
                if (du >= INF_VAL) continue;
                int start = g_row_offsets[u];
                int end   = g_row_offsets[u + 1];
                for (int idx = start; idx < end; idx++) {
                    float nd = du + g_weights[idx];
                    if (nd < g_dist_buf[g_col_indices[idx]]) {
                        g_dist_buf[g_col_indices[idx]] = nd;
                        updated = 1;
                    }
                }
            }
            if (!updated) break;
        }

        distances[r] = (tgt >= 0 && tgt < g_V) ? g_dist_buf[tgt] : INF_VAL;
    }
}

// ── Unified compute_only wrapper (auto-migrated) ──
void solution_compute(int V, int E, const int* h_row_offsets, const int* h_col_indices, const float* h_weights, int num_requests, const int* h_sources, const int* h_targets, float* h_distances) {
    _orbench_old_init(V, E, h_row_offsets, h_col_indices, h_weights);
    _orbench_old_compute(num_requests, h_sources, h_targets, h_distances);
}