SanzSkuy: Created page with "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 m..."

2026-06-28T04:04:58Z

Created page with "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 <code>{cols, rows}</code> — the two were swapped. Square m..."

New page

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 <code>{cols, rows}</code> — 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:

<pre>
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
</pre>

Key observations:
* The failing value was exactly <code>0.00</code> — 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 <code>A_rows</code> so that
the per-rank row slice changed shape:

{| class="wikitable"
! <code>A_rows</code> !! 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''' (<code>my_rows ≠ B_cols</code>). Because matmul distributes ''rows''
across ranks while every rank keeps the full <code>B_cols</code>, every
multi-rank run produces non-square slices — so every multi-rank run failed.

=== Step 2 — Misleading kernel source file ===

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

<syntaxhighlight lang="c">
int col = get_global_id(0); // dim 0 = column
int row = get_global_id(1); // dim 1 = row
</syntaxhighlight>

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

<syntaxhighlight lang="c">
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;
</syntaxhighlight>

== Root cause ==

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

<syntaxhighlight lang="cpp">
// WRONG — dimensions swapped relative to the kernel
size_t global[2] = { (size_t)B_cols, (size_t)my_rows };
</syntaxhighlight>

Effect for a non-square slice <code>my_rows = 100</code>, <code>B_cols = 256</code>:

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

When the slice was square (<code>my_rows == B_cols</code>) 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 <code>C</code>. On some shapes none of
the sampled coordinates landed in the zeroed region, so the run reported
<code>verify OK</code> even though part of the output was wrong. This made the
bug look intermittent until the row/column sweep pinned it down.

== Fix ==

In <code>matmul/src/cl_runner.cpp</code>, 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
(<code>if (row >= rowsA || col >= N) return;</code>) discards the extra
work-items.

<syntaxhighlight lang="cpp">
// 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);
</syntaxhighlight>

<code>ts</code> is chosen from <code>CL_KERNEL_WORK_GROUP_SIZE</code> (stored as
<code>max_wg</code> in <code>ClContext</code>) 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 <code>NULL</code> 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:

{| class="wikitable"
! Scenario !! Before !! After
|-
| Phone solo, non-square (128/85/512 rows) || ''FAILED'' || ''OK''
|-
| Cluster <code>A_rows</code> = 300 / 255 (non-square slices) || ''FAILED'' || ''OK''
|-
| <code>./test.sh</code> 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 ===

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

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

=== <code>/tmp</code> is not writable in Termux ===

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

=== Race when stopping the power sampler ===

<code>power_sample.sh</code> only writes its mean on receiving <code>SIGTERM</code>
(via a trap), and the trap runs only after the current <code>sleep</code>
interval finishes. A fixed <code>sleep 0.3; cat</code> after <code>kill</code>
raced that write and read an empty file. Fixes:
* Same SSH session: <code>kill -TERM $p; wait $p</code>, then read.
* Separate SSH sessions (benchmark): poll until the output file is non-empty before reading.

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

== Affected files ==

{| class="wikitable"
! File !! Change
|-
| <code>matmul/src/cl_runner.cpp</code> || Swap global dims to <code>{rows, cols}</code>; explicit local size + round-up; query <code>CL_KERNEL_WORK_GROUP_SIZE</code>
|-
| <code>matmul/include/cl_runner.hpp</code> || Add <code>size_t max_wg</code> to <code>ClContext</code>
|-
| <code>power_sample.sh</code> || Battery rail read via direct <code>cat</code> or <code>su -c</code> fallback
|-
| <code>benchmark.sh</code> || Phone power: write to <code>$ANDROID_HOME</code>, poll output file
|-
| <code>test.sh</code> || Phone power probe: write to <code>$ANDROID_HOME</code>, <code>wait</code> 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 <code>.cl</code> file with the opposite convention sent the first fix in the wrong direction.
* On Termux: <code>/tmp</code> is not writable, and <code>tar</code>-based deploys need a clean rebuild to avoid stale objects.

Verify FAILED in multi-rank runs - Revision history