From 1f3d489df4d0da10628d38591d8cc606b5a465a9 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Fri, 7 Aug 2026 19:11:01 -0700 Subject: [PATCH 1/2] [mis] validate an input graph with more constraints (https://github.com/ORNL/HeCBench/issues/312) readECLgraph only bounded the neighbor list from above. Both operands of that comparison are signed, so a negative entry was accepted and used directly as an index into the nodes-sized nstat allocation, which compute-sanitizer reports as an invalid __global__ read in the findmins kernel. Reject negative entries. Two further crafted inputs make findmins spin forever rather than read out of bounds: a self loop leaves a node unable to beat itself, and an edge stored in only one direction leaves a node waiting on a neighbor that never clears it. Add verifyUndirectedECLgraph for the checks that only hold for undirected graphs and call it from the mis programs. It is deliberately not part of readECLgraph because floydwarshall2 runs on a directed graph. Also compute the allocation and read sizes in size_t so that nodes + 1 cannot overflow, and close the input file on the error paths. Co-authored-by: Cursor --- src/mis-cuda/graph.h | 75 +++++++++++++++++++++++++++++++++++-------- src/mis-cuda/main.cu | 1 + src/mis-hip/main.cu | 1 + src/mis-omp/main.cpp | 1 + src/mis-sycl/main.cpp | 1 + 5 files changed, 66 insertions(+), 13 deletions(-) diff --git a/src/mis-cuda/graph.h b/src/mis-cuda/graph.h index c6ec216d05..bd5808228a 100644 --- a/src/mis-cuda/graph.h +++ b/src/mis-cuda/graph.h @@ -56,7 +56,8 @@ void freeECLgraph(ECLgraph &g) ECLgraph readECLgraph(const char* const fname) { ECLgraph g; - int cnt; + size_t cnt; + size_t nodes1; // g.nodes + 1 computed in size_t to avoid overflowing int int error_status = 0; FILE* f = fopen(fname, "rb"); if (f == NULL) { @@ -78,9 +79,10 @@ ECLgraph readECLgraph(const char* const fname) exit(-1); } - g.nindex = (int*)malloc((g.nodes + 1) * sizeof(g.nindex[0])); - g.nlist = (int*)malloc(g.edges * sizeof(g.nlist[0])); - g.eweight = (int*)malloc(g.edges * sizeof(g.eweight[0])); + nodes1 = (size_t)g.nodes + 1; + g.nindex = (int*)malloc(nodes1 * sizeof(g.nindex[0])); + g.nlist = (int*)malloc((size_t)g.edges * sizeof(g.nlist[0])); + g.eweight = (int*)malloc((size_t)g.edges * sizeof(g.eweight[0])); if ((g.nindex == NULL) || (g.nlist == NULL) || (g.eweight == NULL)) { fprintf(stderr, "ERROR: memory allocation failed\n\n"); error_status = 1; @@ -88,8 +90,8 @@ ECLgraph readECLgraph(const char* const fname) } // check g.nindex - cnt = fread(g.nindex, sizeof(g.nindex[0]), g.nodes + 1, f); - if (cnt != g.nodes + 1) { + cnt = fread(g.nindex, sizeof(g.nindex[0]), nodes1, f); + if (cnt != nodes1) { fprintf(stderr, "ERROR: failed to read neighbor index list\n\n"); error_status = 1; goto release; @@ -113,14 +115,14 @@ ECLgraph readECLgraph(const char* const fname) } // check g.nlist - cnt = fread(g.nlist, sizeof(g.nlist[0]), g.edges, f); - if (cnt != g.edges) { + cnt = fread(g.nlist, sizeof(g.nlist[0]), (size_t)g.edges, f); + if (cnt != (size_t)g.edges) { fprintf(stderr, "ERROR: failed to read neighbor list\n\n"); error_status = 1; goto release; } for (int v = 0; v < g.edges; v++) { - if (g.nlist[v] >= g.nodes) { + if ((g.nlist[v] < 0) || (g.nlist[v] >= g.nodes)) { fprintf(stderr, "ERROR: value in neighbor list must be a valide node index\n"); error_status = 1; goto release; @@ -128,19 +130,18 @@ ECLgraph readECLgraph(const char* const fname) } // check g.eweight (cnt = 0 is fine) - cnt = fread(g.eweight, sizeof(g.eweight[0]), g.edges, f); + cnt = fread(g.eweight, sizeof(g.eweight[0]), (size_t)g.edges, f); if (cnt == 0) { free(g.eweight); g.eweight = NULL; } - else if (cnt != g.edges) { + else if (cnt != (size_t)g.edges) { error_status = 1; fprintf(stderr, "ERROR: failed to read edge weights\n\n"); } - fclose(f); - release: + fclose(f); if (error_status) { freeECLgraph(g); exit(-1); @@ -149,4 +150,52 @@ ECLgraph readECLgraph(const char* const fname) return g; } + +/* Algorithms such as ECL-MIS require a simple undirected graph: a node that is + its own neighbor, or an edge that is stored in only one of its two + directions, makes them spin forever waiting on a node that can never be + resolved. Graphs that are legitimately directed (used by, e.g., + floydwarshall2) must not be passed to this check. */ + +void verifyUndirectedECLgraph(ECLgraph &g) +{ + for (int v = 0; v < g.nodes; v++) { + for (int i = g.nindex[v]; i < g.nindex[v + 1]; i++) { + if (g.nlist[i] == v) { + fprintf(stderr, "ERROR: neighbor list must not contain self loops\n"); + freeECLgraph(g); + exit(-1); + } + if ((i > g.nindex[v]) && (g.nlist[i - 1] >= g.nlist[i])) { + fprintf(stderr, "ERROR: neighbor list of each node must be sorted in increasing order\n"); + freeECLgraph(g); + exit(-1); + } + } + } + + // the sorted neighbor lists allow the reverse edge to be located by bisection + for (int v = 0; v < g.nodes; v++) { + for (int i = g.nindex[v]; i < g.nindex[v + 1]; i++) { + const int u = g.nlist[i]; + int lo = g.nindex[u]; + int hi = g.nindex[u + 1] - 1; + int found = 0; + while (lo <= hi) { + const int mid = lo + (hi - lo) / 2; + if (g.nlist[mid] == v) { + found = 1; + break; + } + if (g.nlist[mid] < v) lo = mid + 1; else hi = mid - 1; + } + if (!found) { + fprintf(stderr, "ERROR: graph must be undirected, edge %d -> %d is not matched by %d -> %d\n", v, u, u, v); + freeECLgraph(g); + exit(-1); + } + } + } +} + #endif diff --git a/src/mis-cuda/main.cu b/src/mis-cuda/main.cu index d97de74dbe..a15d3b4e69 100644 --- a/src/mis-cuda/main.cu +++ b/src/mis-cuda/main.cu @@ -196,6 +196,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("configuration: %d nodes and %d edges (%s)\n", g.nodes, g.edges, argv[1]); printf("average degree: %.2f edges per node\n", 1.0 * g.edges / g.nodes); diff --git a/src/mis-hip/main.cu b/src/mis-hip/main.cu index d7b5cbedff..3cb52875fa 100644 --- a/src/mis-hip/main.cu +++ b/src/mis-hip/main.cu @@ -196,6 +196,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("configuration: %d nodes and %d edges (%s)\n", g.nodes, g.edges, argv[1]); printf("average degree: %.2f edges per node\n", 1.0 * g.edges / g.nodes); diff --git a/src/mis-omp/main.cpp b/src/mis-omp/main.cpp index 21d55d9c39..9225791d44 100644 --- a/src/mis-omp/main.cpp +++ b/src/mis-omp/main.cpp @@ -153,6 +153,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("configuration: %d nodes and %d edges (%s)\n", g.nodes, g.edges, argv[1]); printf("average degree: %.2f edges per node\n", 1.0 * g.edges / g.nodes); diff --git a/src/mis-sycl/main.cpp b/src/mis-sycl/main.cpp index b2c1b95c1d..70692ad976 100644 --- a/src/mis-sycl/main.cpp +++ b/src/mis-sycl/main.cpp @@ -199,6 +199,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("configuration: %d nodes and %d edges (%s)\n", g.nodes, g.edges, argv[1]); printf("average degree: %.2f edges per node\n", 1.0 * g.edges / g.nodes); From 7c3c565d99ac12c3949bed416ac67347c5155d61 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Mon, 10 Aug 2026 08:03:02 -0700 Subject: [PATCH 2/2] [cc, gc] validate undirected graph inputs Reject malformed graph structures before connected-components and graph-coloring kernels run. Co-authored-by: Cursor --- src/cc-cuda/main.cu | 1 + src/cc-hip/main.cu | 1 + src/cc-sycl/main.cpp | 1 + src/gc-cuda/main.cu | 1 + src/gc-hip/main.cu | 1 + src/gc-omp/main.cpp | 1 + src/gc-sycl/main.cpp | 1 + 7 files changed, 7 insertions(+) diff --git a/src/cc-cuda/main.cu b/src/cc-cuda/main.cu index f4d50abef6..8cf0d2e138 100644 --- a/src/cc-cuda/main.cu +++ b/src/cc-cuda/main.cu @@ -385,6 +385,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); const int repeat = atoi(argv[2]); int* nodestatus = NULL; diff --git a/src/cc-hip/main.cu b/src/cc-hip/main.cu index 8c40f6034a..5b43f59764 100644 --- a/src/cc-hip/main.cu +++ b/src/cc-hip/main.cu @@ -376,6 +376,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); const int repeat = atoi(argv[2]); int* nodestatus = NULL; diff --git a/src/cc-sycl/main.cpp b/src/cc-sycl/main.cpp index 5d68e93bee..c8529d3a26 100644 --- a/src/cc-sycl/main.cpp +++ b/src/cc-sycl/main.cpp @@ -467,6 +467,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); const int repeat = atoi(argv[2]); int* nodestatus = (int*) malloc (sizeof(int) * g.nodes); diff --git a/src/gc-cuda/main.cu b/src/gc-cuda/main.cu index e393855a81..0e5fefdf1b 100644 --- a/src/gc-cuda/main.cu +++ b/src/gc-cuda/main.cu @@ -326,6 +326,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("input: %s\n", argv[1]); printf("nodes: %d\n", g.nodes); printf("edges: %d\n", g.edges); diff --git a/src/gc-hip/main.cu b/src/gc-hip/main.cu index f361954d8d..cf04bf7b41 100644 --- a/src/gc-hip/main.cu +++ b/src/gc-hip/main.cu @@ -330,6 +330,7 @@ int main(int argc, char* argv[]) } ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("input: %s\n", argv[1]); printf("nodes: %d\n", g.nodes); printf("edges: %d\n", g.edges); diff --git a/src/gc-omp/main.cpp b/src/gc-omp/main.cpp index 99365da922..fba6835751 100644 --- a/src/gc-omp/main.cpp +++ b/src/gc-omp/main.cpp @@ -309,6 +309,7 @@ int main(int argc, char* argv[]) const int repeat = atoi(argv[3]); ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); printf("input: %s\n", argv[1]); printf("nodes: %d\n", g.nodes); printf("edges: %d\n", g.edges); diff --git a/src/gc-sycl/main.cpp b/src/gc-sycl/main.cpp index 39dd67078d..07269c6bbe 100644 --- a/src/gc-sycl/main.cpp +++ b/src/gc-sycl/main.cpp @@ -345,6 +345,7 @@ int main(int argc, char *argv[]) { printf("input: %s\n", argv[1]); ECLgraph g = readECLgraph(argv[1]); + verifyUndirectedECLgraph(g); const int nodes = g.nodes; const int edges = g.edges;