From 80x to 385x: A Best-Matching-Unit Search at the L2 Roof, Measured Against a Symmetrically Tuned Baseline
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 |
| 64 2 | SparseBin | vec2, tile 2 | 0.56 | 1.5 GiB | 6.4 |
| 128 2 | SparseBin | vec2, tile 2, clustered, C = 2 | 2.06 | 3.4 GiB | 10.1 |
| 256 2 | SparseBin | vec2, tile 4, clustered, C = 4 | 7.60 | 7.2 GiB | 5.6 |
| 512 2 | SparseBin | vec2, tile 4, clustered, C = 8 | 30.2 | 20.5 GiB | 6.2 |
| comparator | edge | comparator s/ep | published s/ep | published margin | tuned s/ep | tuned margin |
|---|---|---|---|---|---|---|
| MedSOM | 32 2 | 37.63 | 2.393 | 15.7 | 0.31 | 121 |
| MedSOM | 64 2 | 197.78 | 3.590 | 55.1 | 0.56 | 353 |
| MedSOM | 128 2 | 792.76 | 9.932 | 79.8 | 2.06 | 385 |
| MedSOM | 256 2 | 3,209.88 | 42.25 | 76.0 | 7.60 | 422 |
| MedSOM | 512 2 | OOM | - | - | 30.2 | - |
| somoclu | 32 2 | 202.9 | 2.393 | 84.8 | 0.31 | 654 |
| 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% |
| 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% |
| lever | status | why |
|---|---|---|
| bmu1-only | scoped out, twice | 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 |
| 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 |
| 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 |
| # | 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 |
| 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 1.1x |
| launch_bounds register caps | No win; spills convert register pressure into cache and memory traffic |
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