Port SpBench's boolean SpGEMM (cuBool/nsparse hash accumulator) as sp… - #21
Open
laesse10 wants to merge 2 commits into
Open
Port SpBench's boolean SpGEMM (cuBool/nsparse hash accumulator) as sp…#21laesse10 wants to merge 2 commits into
laesse10 wants to merge 2 commits into
Conversation
…gemm_hash SpBench (github.com/EgorOrachyov/SpBench) benchmarks boolean sparse `x` and `+` across cuBool, CUSP, cuSPARSE and SuiteSparse. Only cuBool carries its own mathematics -- the rest are library calls -- so the port follows cuBool_MxM down into its vendored nsparse fork, where the product is a five-phase, row-binned hash accumulator rather than any kind of gemm. nsys on a GH200 (300 iterations each of web-Google and roadNet-CA, CUDA 12.9, sm_90) puts 87% of MxM's wall time on the device and splits that device time 50.6% across the two hash passes, 34.2% across the two bin histogram+scatter passes and 13.4% on the row analysis -- no kernel above 23%, so the boundary is the whole pipeline, not its top symbol. The CUB scan/sort leaves are 1.5%. The port keeps the algorithm that ranking describes: product upper bound per row, binning into (0, 32] ... (2048, 4096] tables, symbolic hash count, exclusive scan, re-binning on exact nnz, numeric hash + bitonic sort + compact. Its docstring names the three simplifications (no global-row path, no accumulate-into-C, sequential order within a bin). Validated at S against a set-union oracle and an independent dense-accumulator Gustavson transcription, on rectangular 97x53x131 axes, and -- outside pytest, which has no GPU -- against what cuBool's CUDA backend actually produced for luxembourg_osm and roadNet-CA: identical row pointers and identical column indices, 393,261 and 12,908,450 nonzeros. That comparison needed two fixes to cuBool first; both are upstream bugs, not build breakage, and are reported separately.
…rontend refusal Two things the first commit got wrong. The hash table was declared ONCE and reused by every row. Upstream gives each row its own shared-memory table -- one per thread block, or one 32-slot slice per 4-thread pwarp group -- so rows carry no dependence at all; a single hoisted table invents a loop-carried dependence the library does not have and hands an optimizer a kernel it cannot parallelize without first undoing the port. Declaring it inside the row loop is what upstream's __shared__ actually is. np.empty rather than np.zeros because the next loop overwrites [0, ts) with the sentinel anyway, and nothing ever reads past it; the reference pays 15% for privatisation (400 -> 460 ms at S) instead of 15x. The docstring now says which phases are parallel, which are atomic histograms, and which is a scan, so the next reader does not have to infer it. The kernel also fails tests/test_dace_numeric_agreement.py, and it is not this kernel's defect: simplify's scalar_to_symbol promotes an integer scalar and remove_symbol_indirection then sympifies a name that resolves to a function. Reduced against the pinned dace a4740d4e7 to one loop, one branch and no arrays beyond the arguments; crc16 and dfa are already listed for the same message. Added to NUMERIC_BAD as parse_fail, next to them, with the repro. The parse-only gate (REFUSED) still says ok, so nothing changes there. Verified after the change: numpy, numba (validation SUCCESS), c/cpp/fortran at fp64 AND fp32, jax; the union oracle, the independent Gustavson transcription and the 97x53x131 rectangular case; tree/yaml/levels; every pre-commit hook.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
…gemm_hash
SpBench (github.com/EgorOrachyov/SpBench) benchmarks boolean sparse
xand+across cuBool, CUSP, cuSPARSE and SuiteSparse. Only cuBool carries its own mathematics -- the rest are library calls -- so the port follows cuBool_MxM down into its vendored nsparse fork, where the product is a five-phase, row-binned hash accumulator rather than any kind of gemm.nsys on a GH200 (300 iterations each of web-Google and roadNet-CA, CUDA 12.9, sm_90) puts 87% of MxM's wall time on the device and splits that device time 50.6% across the two hash passes, 34.2% across the two bin histogram+scatter passes and 13.4% on the row analysis -- no kernel above 23%, so the boundary is the whole pipeline, not its top symbol. The CUB scan/sort leaves are 1.5%.
The port keeps the algorithm that ranking describes: product upper bound per row, binning into (0, 32] ... (2048, 4096] tables, symbolic hash count, exclusive scan, re-binning on exact nnz, numeric hash + bitonic sort + compact. Its docstring names the three simplifications (no global-row path, no accumulate-into-C, sequential order within a bin).
Validated at S against a set-union oracle and an independent dense-accumulator Gustavson transcription, on rectangular 97x53x131 axes, and -- outside pytest, which has no GPU -- against what cuBool's CUDA backend actually produced for luxembourg_osm and roadNet-CA: identical row pointers and identical column indices, 393,261 and 12,908,450 nonzeros.
That comparison needed two fixes to cuBool first; both are upstream bugs, not build breakage, and are reported separately.