Jump to content

Verify FAILED in multi-rank runs

From Haidhi Wiki

This page documents a correctness bug in the MPI × OpenCL matmul cluster where multi-rank runs produced zeroed output elements, the investigation that isolated it, and the fix. It also records three secondary issues found along the way.

TL;DR: The OpenCL kernel maps NDRange dimension 0 to the matrix row and dimension 1 to the column, but the host launch code set the global work size as {cols, rows} — the two were swapped. Square matrices hid the bug because the dimensions were equal; any non-square per-rank slice left whole rows/columns with no work-item, so they stayed zero.

Symptom

Single-device runs verified correctly, but every multi-rank run failed the built-in spot-check: 

rank  host          type  device                                    rows   kernel_ms    gflop/s
0     darmogandolz  GPU   NVIDIA GeForce RTX 3050 6GB Laptop GPU      86        0.44      25.41
1     tumbal        CPU   cpu-ivybridge-AMD A8-6410 APU with AMD      85       13.28       0.84
2     localhost     GPU   QUALCOMM Adreno(TM)                         85       10.13       1.10

wall 0.060 s  ·  aggregate 0.55 GFLOP/s  ·  verify FAILED at C[194][106] gpu=0.00 cpu=7256.00

Key observations:

  • The failing value was exactly 0.00 — an output element that was never written, not a rounding error.
  • It happened on every device type (NVIDIA GPU, AMD CPU, Adreno GPU), so it was not device-specific.
  • Single-device square runs (e.g. 256×256×256) always passed.

Investigation

Step 1 — Square vs. non-square

A 3-node cluster run was repeated while varying only A_rows so that the per-rank row slice changed shape:

A_rows Per-rank slice Result
768 256 rows (square: rows = cols) verify OK
300 100 rows (non-square) verify FAILED
255 85 rows (non-square) verify FAILED

The discriminator was clear: the bug appears only when a rank's slice is non-square (my_rows ≠ B_cols). Because matmul distributes rows across ranks while every rank keeps the full B_cols, every multi-rank run produces non-square slices — so every multi-rank run failed.

Step 2 — Misleading kernel source file

The repository still contained matmul/kernels/matmul.cl (marked deleted in git status) whose index convention was: 

int col = get_global_id(0);   // dim 0 = column
int row = get_global_id(1);   // dim 1 = row

An earlier "fix" had been based on this stale file. The kernel actually compiled into the binary was embedded as a string literal in matmul/src/cl_runner.cpp, with the opposite convention: 

int row = get_global_id(0);   // dim 0 = ROW
int col = get_global_id(1);   // dim 1 = COL
if (row >= rowsA || col >= N) return;
C[row * N + col] = acc;

Root cause

The embedded kernel expects the global work size to be {rows, cols}, but the host set it to {cols, rows}: 

// WRONG — dimensions swapped relative to the kernel
size_t global[2] = { (size_t)B_cols, (size_t)my_rows };

Effect for a non-square slice my_rows = 100, B_cols = 256:

  • Dimension 0 ranged over [0, 256) but was used as the row index. The bounds check row >= rowsA discarded rows ≥ 100.
  • Dimension 1 ranged over [0, 100) but was used as the column index. Columns [100, 256) never received a work-item, so those elements of C kept their initial value of zero.

When the slice was square (my_rows == B_cols) the two ranges were identical, so the swap had no visible effect — which is why square runs always passed and the bug hid for so long.

Why verify was sometimes a false negative

The spot-check only samples 8 elements of C. On some shapes none of the sampled coordinates landed in the zeroed region, so the run reported verify OK even though part of the output was wrong. This made the bug look intermittent until the row/column sweep pinned it down.

Fix

In matmul/src/cl_runner.cpp, set the global work size in the order the kernel expects, use an explicit local work-group size, and round each dimension up to a multiple of it. The kernel's existing bounds check (if (row >= rowsA || col >= N) return;) discards the extra work-items. 

// dim 0 = rows, dim 1 = cols — matches the kernel's get_global_id() mapping
size_t ts = cl.max_wg >= 256 ? 16 : (cl.max_wg >= 64 ? 8 : 4);
size_t local[2]  = { ts, ts };
size_t global[2] = {
    ((size_t)(my_rows + ts - 1) / ts) * ts,   // dim 0 = rows
    ((size_t)(B_cols  + ts - 1) / ts) * ts     // dim 1 = cols
};
clEnqueueNDRangeKernel(cl.queue, cl.kernel, 2, nullptr, global, local, 0, nullptr, nullptr);

ts is chosen from CL_KERNEL_WORK_GROUP_SIZE (stored as max_wg in ClContext) so the local size is valid on every device. Rounding the global size up and relying on the bounds check also removes a related hazard: with a NULL local size and an exact global size, some drivers silently drop the remainder work-items when a dimension is not divisible by the driver-chosen local size.

Verification

After the fix (and a clean rebuild on every node), all shapes pass on all three devices, both solo and in the 3-node cluster:

Scenario Before After
Phone solo, non-square (128/85/512 rows) FAILED OK
Cluster A_rows = 300 / 255 (non-square slices) FAILED OK
./test.sh full suite 6/8 9/9

Secondary issues found

These surfaced while reproducing and verifying the fix on real hardware.

Stale object files on the phone

make deploy-android ships source via tar, which preserves modification times. If the phone's build/ already held newer .o files, make skipped recompiling the changed source and the phone kept running old code. Always force a clean rebuild on the phone after a source change: 

ssh -p 8022 <phone> 'cd matmul-src && rm -rf build && \
  cmake -B build && cmake --build build && cp build/matmul ~/matmul'

/tmp is not writable in Termux

The energy/power scripts originally wrote the sampler output to /tmp/.watt. On Android/Termux /tmp is not writable (Permission denied), so the file was empty. Fix: write to the Termux home ($ANDROID_HOME) instead.

Race when stopping the power sampler

power_sample.sh only writes its mean on receiving SIGTERM (via a trap), and the trap runs only after the current sleep interval finishes. A fixed sleep 0.3; cat after kill raced that write and read an empty file. Fixes:

  • Same SSH session: kill -TERM $p; wait $p, then read.
  • Separate SSH sessions (benchmark): poll until the output file is non-empty before reading.

Also note: kill -INT on a backgrounded non-interactive bash script was unreliable in this environment (the trap did not fire while sleep was running); kill -TERM works. The sampler traps both.

Affected files

File Change
matmul/src/cl_runner.cpp Swap global dims to {rows, cols}; explicit local size + round-up; query CL_KERNEL_WORK_GROUP_SIZE
matmul/include/cl_runner.hpp Add size_t max_wg to ClContext
power_sample.sh Battery rail read via direct cat or su -c fallback
benchmark.sh Phone power: write to $ANDROID_HOME, poll output file
test.sh Phone power probe: write to $ANDROID_HOME, wait for sampler exit

Lessons

  • When a GPU kernel writes zeros, suspect work-item coverage (NDRange dimension order / size) before arithmetic.
  • Square test matrices mask dimension-order bugs — always test a non-square shape.
  • An 8-sample spot-check can be a false negative; for a correctness bug, sweep shapes rather than trusting a single pass.
  • Keep the kernel source in one place. A stale, deleted .cl file with the opposite convention sent the first fix in the wrong direction.
  • On Termux: /tmp is not writable, and tar-based deploys need a clean rebuild to avoid stale objects.
Retrieved from "https://wiki.haidhi.my.id/wiki/index.php?title=Verify_FAILED_in_multi-rank_runs&oldid=706"