[mis] validate an input graph with more constraints

[zjin-lcf]

Fixes #312.

Problem

readECLgraph bounded the neighbor list only from above:

if (g.nlist[v] >= g.nodes) { ... }

g.nlist is int* and g.nodes is int, so this is a signed comparison and a negative entry passes it. findmins then uses the value directly in nstat[nlist[i]], where nstat_d is nodes * sizeof(stattype) bytes with stattype being unsigned char, so the read lands below the allocation. Reproduced with the reporter's PoC:

========= Invalid __global__ read of size 1 bytes
=========     at main.cu:77:findmins(int, const int *, const int *, volatile unsigned char *)
=========     Address 0xbddda60402 is out of bounds
=========     and is 687,864,830 bytes before the nearest allocation at 0xbe06a60000 of size 20 bytes

While fixing that I found two more crafted graphs that pass every existing check and make findmins loop forever instead of reading out of bounds. Both were still running when a 20 s timeout killed them, against 0.12 s for internet.egr:

• nodes=2, nindex=[0,1,2], nlist=[0,1] — each node is its own neighbor, so neither nv > nstat[v] nor nv == nstat[v] && v > v holds and the node can never beat itself.
• nodes=2, nindex=[0,1,1], nlist=[1] — the edge is stored in one direction only, so the neighbor never runs the nstat[nlist[i]] = out write that would clear the waiting node's low bit.

Changes

• Reject negative neighbor list entries.
• Add verifyUndirectedECLgraph, which rejects self loops, unsorted or duplicated neighbor lists, and edges missing their reverse direction, and call it from mis-cuda, mis-sycl, mis-hip, and mis-omp.
• Compute allocation and read sizes in size_t so g.nodes + 1 cannot overflow, and fclose the input on the error paths.

verifyUndirectedECLgraph is deliberately not part of readECLgraph. src/mis-cuda/graph.h is shared via -I../mis-cuda with cc-*, gc-*, and floydwarshall2-*, and floydwarshall2's own input CollegeMsg.egr is a directed graph, so requiring symmetry in the reader would break it. The symmetry test uses bisection over each neighbor list, which is why sortedness is checked first.

Testing

Verified on an NVIDIA Tesla M40 with CUDA 12.1, and with the oneAPI DPC++ SYCL compiler targeting the same GPU.

All four bad inputs are now rejected at load time by both mis-cuda and mis-sycl:

mis-memcheck-negative-index.txt  ERROR: value in neighbor list must be a valide node index
mis-memcheck.txt                 ERROR: neighbor index list always starts at value 0
selfloop.egr                     ERROR: neighbor list must not contain self loops
asym.egr                         ERROR: graph must be undirected, edge 0 -> 1 is not matched by 1 -> 0

mis-cuda ./internet.egr 1 still passes under both compute-sanitizer --tool=memcheck and --tool=initcheck with ERROR SUMMARY: 0 errors.

No regressions in the benchmarks sharing the header: cc-cuda on amazon0601.egr reports PASS, gc-cuda colors the same graph, and floydwarshall2-cuda on the directed CollegeMsg.egr still reports results match. I also confirmed that internet.egr, amazon0601.egr, soc-LiveJournal1.egr, and CollegeMsg.egr are all free of self loops and have strictly sorted neighbor lists, so the new checks reject none of the in-tree inputs.

mis-hip and mis-omp were syntax checked only, as this machine has no AMD GPU and the OpenMP offload Makefile needs icpx.

[Copilot]

Pull request overview

Strengthens ECL graph loading and prevents invalid graph structures from causing out-of-bounds access or nontermination.

Changes:

• Validates neighbor indices and safely computes allocation/read sizes.
• Adds undirected, loop-free, sorted adjacency validation.
• Applies validation across MIS, graph-coloring, and connected-components benchmarks.

Reviewed changes

Copilot reviewed 12 out of 12 changed files in this pull request and generated no comments.

Show a summary per file

File	Description
src/mis-cuda/graph.h	Adds robust loading and graph validation.
src/mis-cuda/main.cu	Validates CUDA MIS input.
src/mis-sycl/main.cpp	Validates SYCL MIS input.
src/mis-hip/main.cu	Validates HIP MIS input.
src/mis-omp/main.cpp	Validates OpenMP MIS input.
src/gc-cuda/main.cu	Validates CUDA GC input.
src/gc-sycl/main.cpp	Validates SYCL GC input.
src/gc-hip/main.cu	Validates HIP GC input.
src/gc-omp/main.cpp	Validates OpenMP GC input.
src/cc-cuda/main.cu	Validates CUDA CC input.
src/cc-sycl/main.cpp	Validates SYCL CC input.
src/cc-hip/main.cu	Validates HIP CC input.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.