From 80× to 385×: A Best-Matching-Unit Search at the L2 Roof, Measured Against a Symmetrically Tuned Baseline
arXiv:2609.05138v1 [cs.LG] 4 Sep 2026
Andrew J. Amosa a
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.1× per epoch over the previously published configuration at map sizes from 322 to 5122 , and lifting the margin over the CUDA implementation behind our earlier MEDLINE atlases from ∼80× to ∼385×. The cuSPARSE implementation SparseBin is compared to received every lever with an analogue on its side, and became 2-3× 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.3× – a terminal result rather than a waypoint, and every untested lever was either capped by that bound by construction or measured null. Keywords: self-organizing map, GPU, CUDA, sparse matrix, kernel tuning, roofline
Email address: [email protected] (Andrew J. Amos) URL: https://orcid.org/0000-0002-9145-0212 (Andrew J. Amos)
1. Introduction Every comparison between two machine learning algorithms has a denominator, and it is common to see comparisons that only tune the author’s kernel. A true understanding of the relative merits of two algorithms requires that both be tuned for peak performance. This paper describes how symmetric tuning illuminated the comparison between alternative algorithms for the best-matching-unit (BMU) search that dominates self-organizing map training [6]. At the scale of a MEDLINE corpus - 26.9 million abstracts as sparse binary term vectors - the search is bound by the bandwidth needed to read the codebook every epoch. A companion paper [3] shows that storing the codebook feature-major, with each feature’s weights contiguous, recasts the search as a tiled sparse-dense product in which every loaded weight column is reused across a tile of articles, and that this accelerates the search 4.58.5× at no cost in quality, since an exact-argmin BMU is invariant to how the codebook is stored [3]. That paper compared its novel SOM algorithm against a baseline SOM algorithm constructed with NVIDIA cuSPARSE library tools. Hereafter I will refer to the novel SOM algorithm as SparseBin and the baseline algorithm as cuSPARSE. This article reports the symmetric tuning of both algorithms. Every lever proven on one was offered to the other before any result was reported, structural asymmetries were named rather than left as an unexplained effort gap, predictions were registered before measurement, and nulls were recorded at the same prominence as wins. Making cuSPARSE faster counted as success - and it succeeded. The tuned cuSPARSE reported here is 2-3× faster than the configuration in the original paper [3]. Four levers - tile size, tile-membership clustering, neuron-axis chunking and vectorised loads - accelerated the BMU search 5.6-10.1× per epoch over the previously published configuration, across map sizes from 322 to 5122 . The margin over the CUDA implementation behind our earlier MEDLINE atlases [1, 2] rose from ∼80× to ∼385×. The tuned SparseBin kernel was the first configuration in the sequence to press a hardware bandwidth ceiling rather than sit beneath all of them: it reached 77% of the L2 bandwidth roof while every other unit ran at 40-65%. Section 6 shows that every lever left on the shelf was either capped by that bound by construction or measured null.
2
1.1. Scope • All measures were completed on one RTX 4090. Every optimum here tile size, chunk count, batch, the L2 roof finding itself - is tuned against that card’s 72 MB L2 and register file. The first paper reported a 10242 run on a 141 GB H200, but no tuning was done so it is not mentioned in this article. 1.2. Contributions 1. A measured 5.6-10.1× per-epoch speed-up of the BMU search, and an ∼80× → ∼385× margin over MedSOM (Section 5). 2. A terminal-limiter result: the tuned kernel reaches 77% of the L2 bandwidth roof, bounding further gains at ∼1.3×, with every unexplored lever either capped by that bound by construction or measured null (Section 6). 3. A public artefact that re-derives the headline table from a fresh clone (Section 9). 2. Method: the rules the programme ran under Four rules governed the programme. They were fixed before it began and are stated here as a protocol another comparison could adopt, with what each one cost. The objective was the best algorithm at each map size, not the best configuration of one implementation. Making cuSPARSE faster counted as success. Parity was a rule, not an intention. Every lever proven on one side was offered to the other before either result was reported. Nulls were reported with the same prominence as wins. Section 7 is the register. A lever that returned under 5% at every map size was recorded and abandoned rather than pursued. Predictions were registered before measurement. Section 7 gives every one with its outcome, including the failures. Quality was excluded from the search and restored at the end. Optimising per-epoch time alone made exhaustive search affordable: a converged run costs minutes to hours, a one-epoch timing seconds. Quality was recorded throughout but never used to select, then validated on the winners in a separate phase. 3
3. Background: the design being tuned This section states only what is needed to follow the tuning programme. The layout argument, its proof of best-matching-unit invariance, and the MEDLINE atlas it was built for are in the first paper [3]. A self-organizing map is trained by repeatedly finding, for each article, the neuron whose weight vector is closest to it - the best-matching unit - and then updating that neuron and its neighbours. On a corpus of sparse binary term vectors the search dominates: it was 98.7-99.5% of epoch time before this programme began. It is bound not by arithmetic but by the bandwidth needed to read the codebook every epoch. The paper’s contribution is that this bound is largely an artefact of how the codebook is stored for a sparse SOM. The SparseBin algorithm stored it feature-major, with each feature’s weights contiguous at W[v.M+i] for feature v and neuron i, which recasts the search as a tiled sparse-dense product in which every loaded weight column is reused across a tile of articles. Because an exact-argmin best-matching unit does not depend on the order weights are held in, the gain is free: held-out quantisation error agrees with cuSPARSE to within 0.5% at every map size. The kernel also fuses the argmin into the product, so the score block is never materialised. cuSPARSE, the comparison algorithm, is built on the proprietary NVIDIA library of that name [7]. It computes the sparse-dense product with a library call, materialises the resulting score block, and then reduces it with a separate argmin pass. Three facts about it matter for the parity rule in section 2. It can select among the library’s algorithms; it can reduce the precision of its score block; and its batch size sets that block’s footprint and therefore its cache behaviour. None of the three has an analogue in a fused kernel, so they are structural asymmetries rather than levers withheld. 4. The levers Four levers produced the speed-up: tile size, tile-membership clustering, neuron-axis chunking, and vectorised loads. They are reported in the order they were run, because the sequence is what identified the binding resource at each stage. A fifth lever, execution ordering, was run and then made redundant by the clustering that followed it; it is kept in the account because it is the experiment that identified the mechanism the clustering went on to exploit. 4
4.1. Tile size The kernel published in the first paper [3] processes the corpus in tiles of 16 articles: for each tile it forms the union of the distinct features those 16 articles touch, then sweeps that set of feature columns once. The tile buys reuse and costs registers - five sixteen-element per-article arrays held live per thread. The tile size was swept across 2, 4, 8, 16 and 32 articles at every map size. Tile 16, the published value, was fastest at none of them. A tile of 2 was 2.7x faster at 1282 , and a tile of 8 was 1.5x faster at 2562 . The mechanism was not the one hypothesised. The expectation was that smaller tiles would cut registers and raise occupancy while raising DRAM traffic, since each loaded feature column is amortised over fewer articles, producing an interior optimum between the two. DRAM traffic turned out to be flat in tile size, within 4% across the whole sweep. The extra codebook re-reads of small tiles are absorbed by L2, which rises from 6.5 to 9.7 TB, and L2 had headroom. The trade is occupancy against L2 traffic, not against DRAM, and occupancy wins by a wide margin. Two further properties emerged and both matter later. The optimum is map-size dependent, so it is a variable rather than a constant. Changing it also perturbs about one assignment in 105 at 2562 , through floating-point tie-breaking rather than any defect in the search. 4.2. Locality: ordering, then membership, then chunking Three levers followed, each aimed at the L2 traffic the tile sweep had identified as binding. Execution ordering. The kernel launches one block per tile of articles in a one-dimensional grid, so the blocks running at any given moment are consecutive tiles in corpus order - which is accession order, and therefore unrelated to what the documents contain. Each block had to read the codebook columns for the features its own articles contained. When two blocks running at the same time need many of the same features, the first block’s reads pull those columns into L2 and the second finds them already there; when their features are disjoint, each block fetches its own columns from DRAM. Permuting which tiles run together, so that simultaneous blocks share as many features as possible, cut DRAM traffic by 43%, 40% and 30% at 642 , 1282 and 2562 at the smallest tile, and raised the L2 hit rate from 49.5% to 64.6% at the largest. Which articles sit in which tile is unchanged, so the 5
order terms are summed in is unchanged, and assignments are bit-identical by construction: twelve of twelve byte comparisons matched. The lever is conditional rather than universal. At a tile of 16 each block already spans sixteen articles, so its own set of features is large and overlaps its neighbours’ heavily without any help; there is little left for reordering to exploit, and at 642 it adds 21% DRAM. Tile membership. Clustering the corpus so that articles sharing rare descriptors fall in the same tile shrinks the union itself rather than merely rescheduling it. This is the stronger form of the same idea, and it made execution ordering a null - the clustered corpus already supplies the locality the ordering was simulating. Ordering is retained in this account as the experiment that identified the mechanism, superseded by its own logic. Neuron-axis chunking. A chunk is a contiguous slice of the map’s neurons. Splitting the neuron axis into C chunks means the kernel sweeps a tile’s feature columns against M/C of the M neurons at a time, taking C passes to cover the whole map, so only one chunk’s weights need be resident at once. The co-resident working set is therefore union x (M/C) x 2 bytes - the tile’s feature union, times the neurons in one chunk, times two bytes per half-precision weight - and is tunable against L2 at any map size. The registered prediction was a residency threshold, a step in the hit rate as the working set crosses the 72 MB cache. The measurement refutes the threshold and confirms the mechanism: the hit rate climbed smoothly from 75.9% to 97.9% across C from 1 to 16, with no step, so chunking buys partial residency and gradual reuse. Occupancy was flat at 83% throughout, which is the diagnostic separating chunking from the tile sweep - two levers that both look like "make the working set smaller" act on different resources, one on cache and one on occupancy. C topped out at 8 at 2562 . C = 16 continued to buy hit rate and cut DRAM to 578 GB, and was slower: each chunk can only find the best match among its own neurons, so the C partial results must be merged into one, and that cost overtakes the memory saving. 4.3. Vectorised loads, and a counter-intuitive optimum Vectorising the codebook loads as __half2 was the largest single lever, 1.37-1.62x. It also moved the chunking optimum from 8 back to 4, which raises DRAM traffic and lowers the hit rate. The optimum is not the configuration that maximises cache statistics; it is the one that balances them against instruction cost. Any tuning account that selects 6
on a cache metric rather than on time would have taken the wrong branch here. 4.4. Symmetrically tuning cuSPARSE and SparseBin Tuning cuSPARSE began as fairness insurance and produced the programme’s most consequential single result. Paper 1’s cuSPARSE comparison spends three quarters of its time at 2562 not in the sparse-dense product but in the argmin read-back that follows it, streaming at 7.7% of measured peak bandwidth. It assigns one thread per article row, so warp-mates stride by M x 2 bytes and every warp access spans 32 sectors: the same uncoalesced access pattern paper 1 identifies as the central error of node-major storage, committed inside my own harness, on cuSPARSE’s side of the comparison. Rewritten as a warp-per-row reduction with vectorised loads it ran 12.5x faster at 2562 , 47.3 s to 3.8 s per epoch, taking DRAM utilisation from 8% to 96%. At smaller maps the gains were 3.5x to 4.3x and the kernel reached 93-96% of peak everywhere. Two further cuSPARSE levers paid. Batch size is the structural analogue of the tile: it sets the score block’s footprint and therefore its cache behaviour, and its optimum is map-size dependent in both directions, with small maps wanting small batches and large maps the largest batch that fits. Algorithm selection is the closest thing to a tile knob a closed library offers, and ALG3 proved a small-map win and a large-map loss when swept alone, and a win at every map size once crossed with batch size on the reordered corpus. That non-additivity is why the programme switched to full factorials: a one-factorat-a-time sweep would have published a cuSPARSE optimum wrong by 1.7x, in SparseBin’s favour, and invisibly. The conclusion this result selects was specified before it was run. The score block is affordable when it is read well, not that fusion was unnecessary. Against a coalesced read-back and a well-chosen batch, fusion confers approximately no per-epoch advantage at 2562 . The advantage of fusion is seen in the memory column of Table 1. 5. Results All timings were taken on one RTX 4090 against the frozen corpus. Section 5.1 reports the comparison the parity rule governs; section 5.2 reports the two comparators it does not. 7
5.1. Per epoch, both implementations tuned One epoch, seed 0, frozen corpus and split, RTX 4090. Median of three replicates, spread ≤1.5% everywhere: Table 1. Per-epoch time with both implementations tuned. One epoch, seed 0, frozen corpus and split, RTX 4090; median of three replicates.
edge
winner
configuration
s/epoch
peak memory
vs published
322
cuSPARSE
0.24
2.5 GiB
3.0×
642 1282
SparseBin SparseBin
0.56 2.06
1.5 GiB 3.4 GiB
6.4× 10.1×
2562
SparseBin
7.60
7.2 GiB
5.6×
5122
SparseBin
warp argmax, ordered, ALG3, 64 MB batch vec2, tile 2 vec2, tile 2, clustered, C = 2 vec2, tile 4, clustered, C = 4 vec2, tile 4, clustered, C = 8
30.2
20.5 GiB
6.2×
5.2. The wider field: MedSOM and somoclu No new campaigns were run - every cell is the frozen record divided by the winner table of Section 5.1. Uniform basis throughout: total wall clock ÷ epochs, both sides. Table 2. Margins over MedSOM and somoclu [10], published and after tuning. Uniform basis: total wall clock divided by epochs, both sides.
comparator
edge
comparator s/ep
published s/ep
published margin
tuned s/ep
tuned margin
MedSOM MedSOM MedSOM MedSOM MedSOM somoclu somoclu somoclu somoclu somoclu
322 642 1282 2562 5122 322 642 1282 2562 5122
37.63 197.78 792.76 3,209.88 OOM 202.9 1,251.5 6,168.5 27,685.0 not run
2.393 3.590 9.932 42.25 2.393 3.590 9.932 42.25 -
15.7× 55.1× 79.8× 76.0× 84.8× 348.6× 621.1× 655.3× -
0.31 0.56 2.06 7.60 30.2 0.31 0.56 2.06 7.60 30.2
121× 353× 385× 422× 654× 2,235× 2,994× 3,643× -
8
speed-up factor (log scale)
Matched work is the first paper’s Section 5.5 definition: a fixed 20 epochs over the full corpus (26,912,934 rows per epoch), the same 90/10 split, the same map, total wall ÷ epochs. The tuned denominators are oneepoch runs, three replicates, spread ≤1.5%, with KL instrumentation left in - so they are conservative. 1282 is quoted throughout the first paper, and it is not the best edge. At 2562 the margins are larger, 422× and 3,643×; quoting 1282 keeps continuity with the first paper’s published figures rather than selecting the maximum. Neither MedSOM nor somoclu was retuned. The parity rule was scoped to the two sparse implementations that compete on the same design, cuSPARSE and SparseBin; MedSOM and somoclu ran at their published or default configurations, so these two margins rest on weaker ground than those in Table 1.
103
79.8× 385×
102 MedSOM, published MedSOM, tuned
32²
64²
map size
somoclu, published somoclu, tuned
128²
256²
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.
5.2.1. A basis inconsistency in paper 1, which this paper does not propagate Paper 1’s abstract and Section 5.5 state ∼82× against MedSOM and 621× against somoclu in the same sentence. Those two figures do not share 9
a denominator. The 621× divides by the total-wall per-epoch figure of 9.932 s; the 82× divides by a training-loop figure of ∼9.67 s that excludes initialisation and evaluation amortisation. On the uniform basis used in the table above, the MedSOM margin is 79.8×, a discrepancy of under 3%. The title rounds that to 80×. This paper prints the recomputable figure. Where paper 1’s number is referred to, it is attributed - "82× as published" - rather than reproduced as though it were derived here. A paper arguing that comparisons should be checkable cannot carry a headline figure that does not recompute from its own artefact. This is also a correction item for paper 1’s journal version. It changes no conclusion; it is a denominator that should be stated once and used twice. 6. The performance ceiling 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.
first addendum (tile 16)
scalar optimum
shipped
registers/thread achieved occupancy DRAM L2
121 33% 45% of peak 0.7 TB/s (14%)
46 66% 41% 3.85 TB/s (77%)
warp-issue busiest pipe limiter
30% ALU 21% latency, under every roof
46 83% 15% 2.5 TB/s (50%) 61% LSU 61% issue/LSU, under every roof
64% LSU 49% at the L2 roof
The shipped kernel was the first configuration in the sequence to press a hardware bandwidth ceiling rather than sit beneath all of them, and the same signature appeared at 5122 . That bounds any future lever at roughly 1.3× - measured as 1.29× whole-epoch at 5122 , 1.19× at 322 .
10
L2 roof reached, 77%
% of peak
80 60 40 20 0
achieved DRAM occupancy first addendum (tile 16)
L2
warp-issue
busiest pipe
scalar optimum
shipped
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.
6.1. Composition of the epoch wall-time The search was 98.7-99.5% of the epoch wall-time when the programme began. Making it roughly six times faster raises the question of whether the rest of the epoch has become material. It has not, and the direction is the opposite of what one would guess. One instrumented epoch at the shipped configuration, free-running clocks, kernel-time sums; the device total matched the trainer’s own wall to within 0.02 s at every map size, so host-side work during the epoch is nil. Table 4. Composition of one instrumented epoch at the shipped configuration, in seconds.
edge
wall
fused BMU
stopping metric
blur
scatter
writeback
norms
BMU share
322 642 1282 2562 5122
0.32 0.56 2.05 7.56 29.98
0.236 0.470 1.817 6.915 28.175
0.033 0.043 0.122 0.358 0.820
0.001 0.003 0.035 0.145 0.592
0.033 0.034 0.037 0.035 0.035
0.001 0.003 0.016 0.064 0.260
0.001 0.002 0.002 0.014 0.065
73.7% 83.9% 88.6% 91.5% 94.0%
11
Under the quantisation-error stop the stopping-metric column disappeared and nothing else moved beyond noise, taking the search’s share to 83.3 / 90.4 / 94.2 / 96.0 / 96.5%. The update’s two components scale as the design predicts, and in opposite directions. The scatter into accumulators is O(N.nnz) and independent of map size: 0.033 to 0.037 s at every rung, across a 256-fold range of neuron counts. The blur is the O(M) term and grows accordingly, 0.001 to 0.592 s. Their sum grows far more slowly than the search does, so the search’s share rises with map size. The single-number epoch was least defensible at the smallest maps, where non-search work is 26% of the epoch, and most defensible at the largest, where it is 3.5%. The whole-epoch ceiling follows. With the search at 96.5% of the epoch at 5122 and the L2 roof bounding search work at about 1.3x, the whole-epoch bound is roughly 1.29x. At 322 , where non-search work is a quarter of the epoch, it is about 1.19x. The update phase costs about one point of the kernel ceiling at large maps and fifteen at small ones, which is the honest bound on any further work of the kind reported here. 6.2. The unpulled levers A number of levers were considered but not pulled. Table 5. The unpulled levers and why each was left.
lever
status
why
bmu1-only
scoped out, twice
shared-memory staging, cp.async nnz binning
scoped out
L1 carveout
measured null
compute/latency lever, reduces no L2 bytes → capped ≲1.3× by construction; and second-best feeds topographic error [5], which both stopping rules gate on at convergence compute/latency levers; neither reduces bytes crossing L2 redistributes instruction work, not L2 traffic the one lever whose mechanism aims at the binding resource - so it was measured rather than argued away
scoped out
Each streaming multiprocessor has a fixed pool of on-chip memory that is divided between the L1 cache and the shared memory a kernel can address, 12
share of epoch wall time (%)
100 80
88.6%
91.5%
94.0%
128² map size
256²
512²
83.9% 73.7%
60 40 20 0
32²
64² fused BMU search stopping metric
blur scatter, writeback, norms
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.
13
and the split can be requested per kernel. Since a larger L1 share absorbs read requests before they reach L2, this is the one untried lever whose mechanism acts on the binding resource itself. The split was swept from the driver’s own choice to both extremes, one epoch per cell, with quantisation error checked identical throughout: Table 6. L1 / shared-memory split sweep on the shipped kernel, seconds per epoch.
edge
driver default
max L1
25% shared
50% shared
max shared
1282 2562 5122
2.07 s 7.55 s 30.11 s
2.12 8.33 32.20
2.24 7.66 30.33
2.24 7.71 30.44
2.24 7.72 30.47
The driver’s own choice was fastest at every map size. Forcing maximum L1, the setting the mechanism argued for, was the worst of the five, by 2.4%, 10.3% and 6.9%. Requesting any fixed split removes the driver’s freedom to vary it, and the kernel uses so little shared memory that the default already favoured L1. The lever is a null, and the fact that it fails on the wrong side of the prediction strengthens the case for leaving it alone. 7. Predictions and nulls 7.1. Registered predictions Seven predictions bearing on the results reported here were registered before the measurement that tested them. Predictions 2 and 6 were the informative failures. Prediction 2 conflated a resident codebook with a resident working set: one stays put, the other turns over, and cross-window eviction eats the margin. Prediction 6 could have been reasoned out in advance - cuSPARSE’s argmin reads the score block, not the codebook, so pinning hot codebook columns was never going to reach its bottleneck. 7.2. Nulls and refutations 8. The cost of symmetry The cost of symmetrically tuning both algorithms was roughly 11.5 GPUhours and 16 hours of implementation on this side; 4.6 and 9 on cuSPARSE, for the tuning programme alone. The validation, protocol-correction and 14
Table 7. Predictions registered before the measurements that tested them. #
Registered
Outcome
1
Ordering reduces L2 traffic >=20% at tile 2, 1282
2
At 1282 , ordering makes the working set resident: hit >95%, DRAM down >2x At 2562 , ordering moves the tile optimum down and DRAM below saturation Chunking tunes residency, with a threshold as the working set crosses L2
Miss on L2 bytes (-0.8%); hit on DRAM (-40%) under the metric reframed before measurement Miss - 82.4%, 1.66x
3 4
5 6 7
The 2562 dead-unit anomaly does not reproduce at the new optimum Hot-column persistence helps cuSPARSE materially more Ordering causes cuSPARSE’s epoch inflation
Hit Split - mechanism confirmed, threshold refuted; smooth partial residency Hit Refuted - helped neither Refuted - one epoch; the cause was a protocol error in that campaign
packaging campaigns that followed added about another 18 GPU-hours, split roughly 55:45 across the two sides. The gap is real and is not a fairness failure. cuSPARSE’s core is a closed library, so its one tunable kernel received its retune and had no vectorisation left to receive, while SparseBin has more of its own code to tune because it is its own code. That a bespoke kernel remains tunable where a library call does not is part of the finding rather than an apology for it. Every lever with an analogue on both sides was run on both. The ratio matters more than the total: the parity work was about a third of the effort, and most of that third was spent once, on understanding cuSPARSE well enough to know what to offer it. 9. Availability of code and data Everything needed to repeat the measurements in this paper is public. The repository sparsesom-tuning holds the tuning and validation campaign: the scripts exactly as they were run, the per-phase timings behind every table above, the profiler captures [8] they were derived from, and a record of which binary and source revision produced each campaign. The corpus itself is deposited separately and openly, so a reader can start from the same 29.9 15
Table 8. Nulls and refutations. 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 Null on both sides, under 1%
Hot-column L2 persistence, both implementations Frequency-ordered feature relabel Reduced-precision score block CSC tile transpose
__launch_bounds__ register caps Second-best-unit tracking removed Chunking at C = 16 Armed Kaski-Lagus evaluation
Execution ordering CUSPARSE_SPMM_CSR_ALG1
Timing null, <=1.5%; retained as substrate Null by construction - already half precision in every published comparison Retired by measurement: its target fell from 19% to 10% of the instruction stream when the tile narrowed. Ceiling ∼1.1x No win; spills convert register pressure into cache and memory traffic Correct lever, wrong bottleneck at the published tile; superseded Fits with a two-buffer merge and buys nothing Degenerate. The trigger fired at epoch 3 in every run and never disarmed, because the improvement series is non-monotone; no choice of threshold fixes it Superseded by clustered membership, which supplies the same locality 4-6x slower than the default everywhere. A hazard, recorded as a warning
16
million documents rather than approximating them. The repository is sparsesom-tuning at release v2.0, archived at Zenodo concept DOI 10.5281/zenodo.22245712. A SHA-256 manifest covers every file, and the record marks which campaigns are the corrected runs. One campaign was withdrawn during the programme after a protocol error was found in it; it is retained in the repository and labelled as withdrawn rather than deleted. The repository was tested as a stranger would use it before release: cloned fresh into a clean directory, built from the committed scripts alone, and used to re-derive the headline table from the corpus at its public DOI. Seven of ten per-epoch cells reproduced exactly, and three carry each binary’s first-launch just-in-time compilation, documented with warm-up guidance. That exercise earned its place twice. The first attempt failed on a virgin configure, where the default host compiler’s standard-library types break the CUDA front end - invisible in a campaign whose build trees were configured months earlier. Both defects it found were packaging defects and no measurement moved, which is the outcome one hopes for and cannot assume. 10. Related work 10.1. Symmetrical tuning Automated tuning. Search-based autotuners - ATLAS in the denselinear-algebra tradition, and general frameworks such as OpenTuner, CLTune and Kernel Tuner - exist precisely to relieve an author of hand-tuning, and an obvious question is why this programme was run by hand. Two reasons, both visible in the results. First, the levers interact non-monotonically: vectorisation moved the chunking optimum from 8 back to 4, and moved it in the direction that raises DRAM traffic and lowers hit rate (Section 4). A search that optimises one parameter at a time converges to a local optimum and reports it as an answer, and the parameter that has to move to escape it is not the one being searched. Second, and more to the point of this paper, an autotuner tunes the parameters you expose to it. On cuSPARSE’s side those are the library’s algorithm selector, its batch size, and the precision of its score block - which is a real search space, and was searched, but it is not the same act as asking what cuSPARSE’s bottleneck actually is and finding that its argmin read-back was uncoalesced. Roofline and limiter analysis. The ceiling argument in Section 6 rests on the roofline model [9], and specifically on its instruction-roofline 17
formulation for GPUs [4], which the companion profiling addendum uses directly. The contribution here is not methodological: it is that a tuning programme was run until the kernel reached a roof, and that the roof is then used to bound what remains rather than to characterise what was achieved. Self-organizing maps at scale. Kohonen’s formulation [6], and the parallel implementations that followed it - somoclu as the multicore-CPU reference [10], and MedSOM as the CUDA implementation behind our earlier atlases [1, 2]. This is context rather than contest; Section 5.3 reports the margins and states plainly which of these the parity rule was applied to. 11. Conclusion Tuning the best-matching-unit search through four standard levers made it 5.6-10.1× faster per epoch than the configuration previously published. The comparison it is set against is not the one the earlier paper used: cuSPARSE received every applicable lever from the same programme and became 2-3× faster in the process. The margin survives that improvement, which is the only reason it is worth reporting. The search stopped at a ceiling rather than at exhaustion. The tuned kernel pressed the L2 bandwidth roof at 77% of peak while every other unit sat at 40-65%, which bounds any further lever on this device at roughly 1.3× - measured as 1.29× whole-epoch at 5122 and 1.19× at 322 . Every lever left untried is accounted for against that bound, and the one whose mechanism aimed at the roof itself was measured rather than argued away. This is a single-device study. The values here are specific to one consumer card, and I expect the balance to move on a part with different cache and bandwidth proportions. What transfers is not the tile size or the chunk count. It is the observation that a comparison’s denominator is a variable, that it is worth two to three times in this case, and that reporting it costs about a third of the effort already being spent. References [1] Amos, A. J., Lee, K., Sen Gupta, T. & Malau-Aduli, B. S. (2024a). Defining the boundaries of psychiatric and medical knowledge: applying cartographic principles to self-organising maps. In MEDINFO 2023 – The Future Is Accessible, Studies in Health Technology and Informatics, vol. 310, pp. 795–799. Amsterdam: IOS Press. doi:10.3233/SHTI231074 18
[2] Amos, A. J., Lee, K., Sen Gupta, T. & Malau-Aduli, B. S. (2024b). Validating the knowledge represented by a self-organizing map with an expert-derived knowledge structure. BMC Medical Education, 24, 416. doi:10.1186/s12909-024-05352-y [3] Amos, A. J. (2026). A feature-major codebook for memory-efficient sparse-binary self-organizing maps: scaling a MEDLINE atlas to 1.05 million neurons on a single consumer GPU. arXiv:2608.24067 [cs.LG]. doi:10.48550/arXiv.2608.24067 [4] Ding, N. & Williams, S. (2019). An instruction roofline model for GPUs. In 2019 IEEE/ACM Performance Modeling, Benchmarking and Simulation of High Performance Computer Systems (PMBS), pp. 7–18. doi:10.1109/PMBS49563.2019.00007 [5] Kaski, S. & Lagus, K. (1996). Comparing self-organizing maps. In Artificial Neural Networks – ICANN’96, Lecture Notes in Computer Science 1112, pp. 809–814. Berlin: Springer. doi:10.1007/3-540-61510-5_136 [6] Kohonen, T. (2013). Essentials of the self-organizing map. Neural Networks, 37, 52–65. doi:10.1016/j.neunet.2012.09.018 [7] NVIDIA Corporation (2025). cuSPARSE Library, CUDA Toolkit 12.8. https://docs.nvidia.com/cuda/cusparse/ [accessed 2026-07-30]. [8] NVIDIA Corporation (2025). Nsight Compute, CUDA Toolkit 12.8. https://docs.nvidia.com/nsight-compute/ [accessed 2026-07-30]. [9] Williams, S., Waterman, A. & Patterson, D. (2009). Roofline: an insightful visual performance model for multicore architectures. Communications of the ACM, 52(4), 65–76. doi:10.1145/1498765.1498785 [10] Wittek, P., Gao, S. C., Lim, I. S. & Zhao, L. (2017). somoclu: An efficient parallel library for self-organizing maps. Journal of Statistical Software, 78(9), 1–21. doi:10.18637/jss.v078.i09
19