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

Andrew James Amos

cs.LGSep 4, 2026

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

Authors: Andrew James Amos

Organizations: College of Medicine and Dentistry, James Cook University, Townsville, Queensland, Australia

Abstract

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.

Figures & tables

edge	winner	configuration	s/epoch	peak memory	vs published
32 2	cuSPARSE	warp argmax, ordered, ALG3, 64 MB batch	0.24	2.5 GiB	3.0 $\times$
64 2	SparseBin	vec2, tile 2	0.56	1.5 GiB	6.4 $\times$
128 2	SparseBin	vec2, tile 2, clustered, C = 2	2.06	3.4 GiB	10.1 $\times$
256 2	SparseBin	vec2, tile 4, clustered, C = 4	7.60	7.2 GiB	5.6 $\times$
512 2	SparseBin	vec2, tile 4, clustered, C = 8	30.2	20.5 GiB	6.2 $\times$

Table 1: Per-epoch time with both implementations tuned. One epoch, seed 0, frozen corpus and split, RTX 4090; median of three replicates.

comparator	edge	comparator s/ep	published s/ep	published margin	tuned s/ep	tuned margin
MedSOM	32 2	37.63	2.393	15.7 $\times$	0.31	121 $\times$
MedSOM	64 2	197.78	3.590	55.1 $\times$	0.56	353 $\times$
MedSOM	128 2	792.76	9.932	79.8 $\times$	2.06	385 $\times$
MedSOM	256 2	3,209.88	42.25	76.0 $\times$	7.60	422 $\times$
MedSOM	512 2	OOM	-	-	30.2	-
somoclu	32 2	202.9	2.393	84.8 $\times$	0.31	654 $\times$

Table 2: Margins over MedSOM and somoclu [ 10 ] , published and after tuning. Uniform basis: total wall clock divided by epochs, both sides.

Figure 1: Speed-up over MedSOM and somoclu, as published and after tuning, by map size (log scale). The MedSOM pair is a GPU-to-GPU comparison; the somoclu pair is against a multicore-CPU library. Neither comparator was retuned.

	first addendum (tile 16)	scalar optimum	shipped
registers/thread	121	46	46
achieved occupancy	33%	83%	66%
DRAM	45% of peak	15%	41%
L2	0.7 TB/s (14%)	2.5 TB/s (50%)	3.85 TB/s (77%)
warp-issue	30%	61%	64%
busiest pipe	ALU 21%	LSU 61%	LSU 49%

Table 3: Limiter progression across the programme. The shipped kernel is the first configuration to press a hardware roof rather than sit beneath all of them.

Figure 2: Limiter progression across the programme. The shipped kernel is the first configuration to press a hardware ceiling rather than sit beneath all of them.

edge	wall	fused BMU	stopping metric	blur	scatter	write- back	norms	BMU share
32 2	0.32	0.236	0.033	0.001	0.033	0.001	0.001	73.7%
64 2	0.56	0.470	0.043	0.003	0.034	0.003	0.002	83.9%
128 2	2.05	1.817	0.122	0.035	0.037	0.016	0.002	88.6%
256 2	7.56	6.915	0.358	0.145	0.035	0.064	0.014	91.5%
512 2	29.98	28.175	0.820	0.592	0.035	0.260	0.065	94.0%

Table 4: Composition of one instrumented epoch at the shipped configuration, in seconds.

Figure 3: Composition of one instrumented epoch at the shipped configuration. The search’s share rises with map size, so the single-number epoch is least defensible at the smallest maps and most defensible at the largest.

lever	status	why
bmu1-only	scoped out, twice	compute/latency lever, reduces no L2 bytes $\rightarrow$ capped $\lesssim$ 1.3 $\times$ by construction; and second-best feeds topographic error [ 5 ] , which both stopping rules gate on at convergence
shared-memory staging, cp.async	scoped out	compute/latency levers; neither reduces bytes crossing L2
nnz binning	scoped out	redistributes instruction work, not L2 traffic
L1 carveout	measured null	the one lever whose mechanism aims at the binding resource - so it was measured rather than argued away

Table 5: The unpulled levers and why each was left.

edge	driver default	max L1	25% shared	50% shared	max shared
128 2	2.07 s	2.12	2.24	2.24	2.24
256 2	7.55 s	8.33	7.66	7.71	7.72
512 2	30.11 s	32.20	30.33	30.44	30.47

Table 6: L1 / shared-memory split sweep on the shipped kernel, seconds per epoch.

#	Registered	Outcome
1	Ordering reduces L2 traffic >=20% at tile 2, 128 2	Miss on L2 bytes (-0.8%); hit on DRAM (-40%) under the metric reframed before measurement
2	At 128 2 , ordering makes the working set resident: hit >95%, DRAM down >2x	Miss - 82.4%, 1.66x
3	At 256 2 , ordering moves the tile optimum down and DRAM below saturation	Hit
4	Chunking tunes residency, with a threshold as the working set crosses L2	Split - mechanism confirmed, threshold refuted; smooth partial residency
5	The 256 2 dead-unit anomaly does not reproduce at the new optimum	Hit
6	Hot-column persistence helps cuSPARSE materially more	Refuted - helped neither

Table 7: Predictions registered before the measurements that tested them.

Lever	Outcome
cusparseSpMM_preprocess and persistent descriptors	Null at every map size and algorithm. The defect was real; its cost was invisible at one-epoch granularity
Hot-column L2 persistence, both implementations	Null on both sides, under 1%
Frequency-ordered feature relabel	Timing null, <=1.5%; retained as substrate
Reduced-precision score block	Null by construction - already half precision in every published comparison
CSC tile transpose	Retired by measurement: its target fell from 19% to 10% of the instruction stream when the tile narrowed. Ceiling $\sim$ 1.1x
launch_bounds register caps	No win; spills convert register pressure into cache and memory traffic

Table 8: Nulls and refutations.

GPU Kernel Optimization

GPU kernel optimization improves the speed and memory efficiency of computations by tailoring code to GPU hardware and workloads.

See all 210 papers

FloatSOM: GPU-Accelerated, Distributed, Topology-Flexible Self-Organizing Maps
Apr 29, 2026Tony Xu, Sarah Klamt, Katherine Turner +3GPU Acceleration High-Performance Computing
John Curtin School of Medical Research, Australian National University · Department of Information Technology and Electrical Engineering, ETH Zürich · Mathematical Data Science Centre, Australian National University +2
SparseDitto: An Agentic Sparse Compilation Framework through Architecture-Aware Synthesis on GPUs
Aug 5, 2026Shiyang Li, Guangyan Sun, Jinwei Tang +3GPU Kernel Optimization High-Performance Computing
University of Minnesota Minneapolis, MN, USA · Northeastern University Boston, MA, USA
Fast LapSum: Exact Differentiable Top- $k$ at Million Scale
Aug 7, 2026Jakub Antczak, Joanna Wojciechowicz, Kamil Książek +3Differentiable Programming Differentiable Optimization
Wrocław University of Science and Technology Wrocław, Poland · Centre for Credible Artificial Intelligence Warsaw University of Technology Warsaw, Poland · Jagiellonian University Kraków, Poland

Apr 29, 2026J/K move · Enter open · S save

FloatSOM: GPU-Accelerated, Distributed, Topology-Flexible Self-Organizing Maps

Tony Xu, Sarah Klamt, Katherine Turner, Anne Brustle+2

John Curtin School of Medical Research, Australian National University · Department of Information Technology and Electrical Engineering, ETH Zürich · Mathematical Data Science Centre, Australian National University+2

GPU Acceleration High-Performance Computing Distributed Training Self-Organizing Maps

Open paper AskPDF arXiv

GPU-accelerated Self-Organizing Map (SOM) implementations are among the most competitive options for large-scale SOM analysis, but growing dataset sizes increasingly challenge their practical use because workloads no longer fit cleanly within device-memory limits. We introduce FloatSOM, a SOM framework for scalable training and deployment that supports multi-GPU execution, out-of-memory disk-backed streaming, and novel topologies beyond regular lattices. We evaluate FloatSOM on 14 synthetic and real benchmark datasets together with controlled speed scaling benchmarks, and show that these improved topologies, combined with topology-aware hyperparameter fine-tuning, yield lower quantization error than current state-of-the-art SOM baselines. FloatSOM also sustains this performance at large scale with high-throughput distributed execution; in the largest benchmark, it trains a 1024-node SOM network on 1,000,000,000 samples with 50 features in 6.16 minutes on 8 GPUs across two separate high-performance-computing nodes.

Abstract

GPU Kernel Optimization

Explore similar work

FloatSOM: GPU-Accelerated, Distributed, Topology-Flexible Self-Organizing Maps

SparseDitto: An Agentic Sparse Compilation Framework through Architecture-Aware Synthesis on GPUs

Fast LapSum: Exact Differentiable Top-kkk at Million Scale

Fast LapSum: Exact Differentiable Top- $k$ at Million Scale