PaperlayerView exact revision on arXiv

Source-linked AI summary

From 80x to 385x: A Best-Matching-Unit Search at the L2 Roof, Measured Against a Symmetrically Tuned Baseline

Andrew James Amos

arXiv:2609.05138v1cs.LG

TL;DR

GPU algorithm comparisons can be misleading when only one implementation is tuned. This paper symmetrically tunes SparseBin and cuSPARSE for BMU search, finding a substantially larger surviving margin and a bandwidth-bound stopping point.

  • Problem

    GPU algorithm comparisons often tune only the author’s kernel, limiting fair assessment of relative performance.

  • Method

    The study symmetrically tunes SparseBin and cuSPARSE using four BMU-search levers: tile size, tile-membership clustering, neuron-axis chunking, and vectorised loads.

  • Results

    SparseBin became 5.6-10.1× faster per epoch, while the margin over the CUDA implementation rose from ∼80× to ∼385× after cuSPARSE improved 2-3×.

  • Takeaways & Limitations

    The comparison remains strongly favorable to SparseBin after symmetric tuning, while the tuned kernel reaches a hardware bandwidth ceiling that bounds further gains at roughly 1.3×.

  • Takeaways & Limitations

    The tile-size optimum is map-size dependent, and changing it can perturb assignments through floating-point tie-breaking.

Abstract

from arXiv · show

Comparisons between GPU implementations are usually asymmetric: one side is tuned by its author, the other is run as found. I report a programme that tuned both a novel SOM algorithm (SparseBin) and the baseline algorithm it was being compared to (cuSPARSE). The best-matching-unit search that dominates self-organizing map training was tuned through four levers - tile size, tile-membership clustering, neuron-axis chunking and vectorised loads - reaching 5.6-10.1x per epoch over the previously published configuration at map sizes from 32x32 to 512x512, and lifting the margin over the CUDA implementation behind our earlier MEDLINE atlases from ~80x to ~385x. cuSPARSE, the implementation SparseBin is compared against, received every lever with an analogue on its side, and became 2-3x faster in the process. The tuned kernel pressed the L2 bandwidth roof at 77% of peak with every other unit at 40-65%, bounding any further lever at ~1.3x - a terminal result rather than a waypoint, and every untested lever was either capped by that bound by construction or measured null.

1. Introduction

The paper applies symmetric tuning to SparseBin and cuSPARSE, making the comparison fairer while substantially improving both implementations. Four levers raise BMU-search speed and expose an L2-bandwidth limit on further gains.

  • Fair comparison: Symmetric tuning gave cuSPARSE every applicable analogue before results were reported, making its 2-3× improvement part of the comparison.The tuned baseline was not left at its original published configuration.
  • Performance result: 5.6-10.1× per epoch was achieved over the previously published configuration across map sizes from 32×32 to 512×512.The four levers were tile size, tile-membership clustering, neuron-axis chunking, and vectorised loads.
  • Performance result: ∼80× → ∼385× was the resulting margin over the CUDA implementation used for earlier MEDLINE atlases.This margin is reported after improving the comparison baseline.
  • Roofline result: 77% of the L2 bandwidth roof was reached, with every other unit at 40-65%, bounding further gains at roughly 1.3×.The paper treats this as a terminal bound; remaining levers were either capped by it or measured null.
  • Scope: All measurements and optima were tuned on one RTX 4090, including its 72 MB L2 and register file.The supplied scope passage notes that an H200 run from the first paper was not tuned here.

2. Method: the rules the programme ran under

The tuning programme defined fairness as a protocol rather than an intention: both implementations received applicable levers, predictions were registered, and null results were retained. Timing drove configuration selection, while quality was validated separately on winners.

  • Programme objective: The objective was the best algorithm at each map size, not the best configuration of one implementation.Making cuSPARSE faster counted as success under the programme’s objective.
  • Parity: Every lever proven on one side was offered to the other before either result was reported.This made parity an explicit rule governing the comparison.
  • Null handling: Levers returning under 5% at every map size were recorded and abandoned rather than pursued.Nulls received the same prominence as wins.
  • Measurement protocol: Predictions were registered before measurement, including failures.The programme kept a pre-registered record of predicted outcomes and measured results.
  • Selection and validation: Per-epoch time selected configurations, while quality was recorded throughout and validated separately on the winners.This separated inexpensive timing runs from converged quality evaluation.

3. Background: the design being tuned

The tuned design targets the bandwidth-dominated BMU search in sparse SOM training. SparseBin reorganizes the codebook and fuses reduction with the product, while cuSPARSE retains library-specific structural asymmetries.

  • Workload: The BMU search dominated 98.7-99.5% of epoch time before tuning and was limited by codebook bandwidth rather than arithmetic.Each article’s closest neuron is found before updating that neuron and its neighbours.
  • SparseBin design: SparseBin stores the codebook feature-major, reusing loaded weight columns across article tiles in a sparse-dense product.The exact-argmin BMU is invariant to weight storage order, and held-out quantisation error agrees with cuSPARSE within 0.5%.
  • SparseBin design: SparseBin fuses the argmin into the product so the score block is never materialised.This is part of the design being tuned rather than a separate post-processing pass.
  • cuSPARSE design: cuSPARSE materialises the sparse-dense score block and reduces it with a separate argmin pass.Its library algorithm choice, score precision, and batch size affect cache behaviour and have no direct fused-kernel analogue.

4. The levers

The programme tuned SparseBin and cuSPARSE through interacting levers that target occupancy, cache locality, working-set size, and memory access. These experiments identified map-size-dependent optima, conditional gains, and a bandwidth-bound ceiling.

  • Overview: Four levers—tile size, tile-membership clustering, neuron-axis chunking, and vectorised loads—were tuned in sequence to identify the binding resource at each stage.Execution ordering was also tested, then made redundant by clustering.
  • 4.1. Tile size: Tile 16, the published value, was fastest at none of the tested map sizes; tile 2 was 2.7x faster at 128^2, while tile 8 was 1.5x faster at 256^2.The tile sweep covered 2, 4, 8, 16, and 32 articles.
  • 4.1. Tile size: Smaller tiles traded occupancy against L2 traffic rather than DRAM traffic, because extra codebook rereads were absorbed by L2.DRAM traffic stayed within 4% across the sweep, while L2 rose from 6.5 to 9.7 TB.
  • 4.2. Locality: ordering, then membership, then chunking: 43%, 40%, and 30% DRAM reductions at 64^2, 128^2, and 256^2 followed from scheduling tiles with shared features, while assignments remained bit-identical.At the largest map, the L2 hit rate rose from 49.5% to 64.6%; the benefit was conditional and could add 21% DRAM at 64^2 with tile 16.
  • 4.2. Locality: ordering, then membership, then chunking: Neuron-axis chunking raised the L2 hit rate smoothly from 75.9% to 97.9% as C increased from 1 to 16, showing partial residency rather than a threshold.At 256^2, C = 16 was slower despite lower DRAM because partial results required merging.
  • 4.3. Vectorised loads, and a counter-intuitive optimum: 1.37-1.62x came from vectorised __half2 loads, but they moved the chunking optimum from 8 back to 4, demonstrating that cache metrics alone can select the wrong configuration.The optimum balanced cache behaviour against instruction cost.
  • 4.4. Symmetrically tuning cuSPARSE and SparseBin: 12.5x faster cuSPARSE read-back at 256^2 reduced epoch time from 47.3 s to 3.8 s and raised DRAM utilisation from 8% to 96%.The rewrite used a warp-per-row reduction with vectorised loads; batch size and algorithm selection were then tuned jointly because their effects were non-additive.
  • 4.4. Symmetrically tuning cuSPARSE and SparseBin: Against coalesced read-back and a well-chosen batch, fusion conferred approximately no per-epoch advantage at 256^2.The result supports comparing complete tuned pipelines rather than attributing performance to fusion alone.

5. Results

Results are reported on a uniform wall-clock-per-epoch basis using frozen data and repeated measurements, with both sparse implementations tuned. The recalculated comparison preserves the reported conclusion while distinguishing recomputable margins from inherited published figures and noting that MedSOM and somoclu were not retuned.

  • Median timings used one epoch on an RTX 4090 with the frozen corpus and split, with three replicates and spread ≤1.5%.
  • Per-epoch comparisons divide total wall clock by epochs for both implementations.
  • At 256^2, the margins over MedSOM and somoclu are 422× and 3,643×, respectively.The paper quotes 128^2 for continuity with earlier published figures rather than selecting the maximum edge.
  • 79.8× is the uniform-basis MedSOM margin, compared with 82× as published, because the earlier figures used different denominators.The 621× somoclu figure divides by a total-wall per-epoch value, while the 82× MedSOM figure used a training-loop value excluding initialization and evaluation amortization.
  • The headline table is recomputable from the paper’s artefact, but MedSOM and somoclu were left at published or default configurations rather than retuned.The parity rule applied only to SparseBin and cuSPARSE, so these wider margins rest on weaker ground than the symmetric sparse-implementation comparison.

6. The performance ceiling

The shipped kernel is the first configuration to press a hardware roof, with search performance approaching an L2 bandwidth ceiling. The remaining epoch work and tested unpulled levers leave only limited further improvement.

  • Composition of the epoch wall-time: The search occupied 98.7-99.5% of epoch wall-time before tuning, and host-side work was nil at the shipped configuration.The device total matched trainer wall-time within 0.02 s at every map size.
  • Composition of the epoch wall-time: The search’s epoch share rose from 83.3% to 96.5% as map size increased because update work grew more slowly.Non-search work was 26% of the epoch at the smallest maps and 3.5% at the largest.
  • The performance ceiling: The L2 roof bounds further search improvement at roughly 1.3×, corresponding to 1.29× whole-epoch improvement at 512² and 1.19× at 32².The update phase accounts for about one point of the kernel ceiling at large maps and fifteen at small ones.
  • The unpulled levers: The L1/shared-memory split was a null: the driver’s choice was fastest at every map size, while maximum L1 was slower by 2.4%, 10.3%, and 6.9%.The default already favored L1 because the kernel uses little shared memory.

7. Predictions and nulls

Seven predictions were registered before measurement, including failures that clarified which assumptions did not hold. The failures distinguished resident codebook data from a changing working set and identified cuSPARSE’s actual bottleneck.

  • Predictions and nulls: Seven predictions bearing on the reported results were registered before measurement.The predictions were tested and their outcomes reported, including failures.
  • Predictions and nulls: Prediction 2 failed because it conflated a resident codebook with a resident working set, whose turnover caused cross-window eviction.Prediction 6 failed because cuSPARSE’s argmin reads the score block rather than the codebook, so hot-codebook pinning could not reach its bottleneck.

8. The cost of symmetry

The comparison explicitly paid for symmetric tuning, offering each algorithm every available analogous lever. The resulting gap remained real despite the effort spent understanding and retuning cuSPARSE.

  • The cost of symmetry: The tuning programme cost roughly 11.5 GPU-hours and 16 implementation hours for SparseBin, versus 4.6 and 9 for cuSPARSE.Validation, protocol-correction, and packaging campaigns added about 18 GPU-hours.
  • The cost of symmetry: Every lever with an analogue on both sides was run on both, while cuSPARSE’s structural asymmetries were named rather than treated as withheld effort.cuSPARSE’s closed library core limited which additional tuning operations existed on its side.
  • The cost of symmetry: Parity work consumed about a third of the effort, mostly to understand cuSPARSE well enough to know what to offer it.Making cuSPARSE faster counted as success, and it became 2-3× faster than the original configuration.

9. Availability of code and data

The paper’s code, measurements, profiler captures, and corpus are publicly available for reproduction. A fresh-clone test reproduced most per-epoch cells exactly and exposed packaging defects without moving measurements.

  • Availability of code and data: The sparsesom-tuning repository contains the executed scripts, per-phase timings, profiler captures, and campaign binary/source records.The corpus is deposited separately and openly, allowing reproduction from the same documents.
  • Availability of code and data: Seven of ten per-epoch cells reproduced exactly from a fresh clone, while three included documented first-launch just-in-time compilation.The repository was rebuilt and used to re-derive the headline table from the public corpus.
  • Availability of code and data: The first clean-configure attempt exposed compiler and packaging defects, but neither defect changed any measurement.The repository retains a withdrawn campaign labelled after a protocol error was found.

10. Related work

The paper situates its contribution in symmetric, roofline-guided GPU tuning rather than automated parameter search alone. It frames self-organizing-map implementations as context for a parity-based comparison.

  • Automated tuning: Vectorisation and chunking interact non-monotonically, so one-parameter-at-a-time autotuning can settle on local optima.The paper reports that vectorisation moved the chunking optimum from 8 back to 4 while increasing DRAM traffic and lowering hit rate.
  • Automated tuning: Autotuners search only the parameters exposed to them, whereas this programme also investigated cuSPARSE’s actual bottleneck: uncoalesced argmin read-back.The cuSPARSE search space included its algorithm selector, batch size, and score-block precision, but not necessarily the bottleneck itself.
  • Roofline and limiter analysis: The roofline analysis is used to bound remaining optimization opportunity after the kernel reaches a hardware ceiling, not merely to characterize achieved performance.The paper identifies this use of the roof as the contribution of the tuning programme rather than a new roofline methodology.
  • Self-organizing maps at scale: Self-organizing maps and implementations such as somoclu and MedSOM provide comparison context, while the paper reports which implementations receive the parity rule.The related-work discussion distinguishes context from the contest itself.

11. Conclusion

The conclusion reports a faster BMU search and a comparison that remains favorable after cuSPARSE receives symmetric tuning. It treats the observed L2 ceiling as the stopping boundary and limits portability claims to the measured device.

  • Conclusion: 5.6-10.1× per epoch: four levers accelerated the BMU search over the previously published configuration.The levers were tile size, tile-membership clustering, neuron-axis chunking, and vectorised loads.
  • Conclusion: 2-3×: symmetric tuning also accelerated cuSPARSE, while the comparison margin survived that improvement.The conclusion presents the surviving margin as the reason for reporting the result.
  • Conclusion: 77% of peak L2 bandwidth: the tuned kernel reached the reported ceiling, bounding further gains on this device at roughly 1.3×.Every untried lever was either capped by that bound by construction or measured null.
  • Conclusion: The study is limited to one consumer GPU, so tile size and chunk count are not claimed to transfer unchanged across devices.The transferable observation is the value of treating the comparison denominator as a variable, not the specific tuned parameters.
Loading 2609.05138v1…