Skip to content

Add launch-box support + Perf tuning - #44

Closed
neoblizz wants to merge 175 commits into
mainfrom
neoblizz/perf
Closed

neoblizz wants to merge 175 commits into
mainfrom
neoblizz/perf

Conversation

@neoblizz

Copy link
Copy Markdown
Member

No description provided.

neoblizz added 30 commits June 17, 2021 12:16
…t it doesn't work when block-size > tile-size, most likely due incorrect pointer arithmetic for shared memory being used. Also doesn't launch the kernel for large datasets like hollywood.
The problem still exists where results are simply incorrect if the block-size is not equal to the tile-size.
We notice that in a small dataset, the result at 0th index is 18 (correct is 11), when at 32nd index it is 0 (correct is 7), which means somehow the 32nd index is accumulating at 0th index (the difference aligns perfectly for all places with errors).

Known problem:
Even with block-size == tile-size, it seems like we have to drop the floating point accuracy to 1e-2 for the result to equal the CPU results. This may be a floating point issue, but also could be related to something else that I am missing.
Note this needs a ton of clean-up to hide the complexity of setting up the loops.
But the beauty of this commit is that block size does not need to equal the tile
size. That means you can have a large block size (meaning, 128 or 256 threads in
a launch per block), and can partition the threads within to several tiles, maybe
ideally equal to the size of a warp (32). And have those work on their own
processing each tile and balancing it. This takes the idea of thread, block or warp
mapped schedule, and turns it into a single group-mapped schedule.
Removed mainly shared memory cluter.
neoblizz and others added 28 commits May 5, 2026 21:25
- Dockerfile: multi-stage (builder + runtime); COPY . /workspace
  instead of git clone gunrock/loops (which baked an upstream
  snapshot ignoring local edits). Pull CMake from Kitware's APT
  repo to stay current. Parameterize CUDA_VERSION and
  UBUNTU_VERSION as ARGs.
- docker-compose.yml: replace devices: /dev/dri + group_add: video
  (DRI render-node convention; AMD/Intel only, never exposed an
  NVIDIA GPU) with the NVIDIA Container Toolkit setup
  (runtime: nvidia + deploy.resources.reservations.devices). Fix
  the dockerfile path: Dockerfile -> docker/Dockerfile.
- Replace three near-duplicate Linux workflows (ubuntu-22.04,
  ubuntu-latest) and the Windows workflow with one build.yml
  matrix: Ubuntu 22.04/24.04 x CUDA 12.6.3/13.0.0 plus Windows
  + CUDA 12.6.3.
- All Linux jobs use the new ci-multi-arch preset (sm_70;sm_80;sm_90).
- Drop LOOPS_BUILD_TESTS=ON (unittests/ is empty; would build no
  test binaries and just slow down CI).
- Remove the windows symlink-deletion hack (no longer needed
  without the standalone Thrust fetch).
- clang-format only commits on push, not PRs from forks.
- Add concurrency: so duplicate runs cancel each other.
Three changes that together unblock Windows CI (which was already
failing on main before this branch):

- Bump Jimver/cuda-toolkit v0.2.27 -> v0.2.35. v0.2.34 was the
  Node 20 -> Node 24 migration; v0.2.27 still ships Node 20 actions
  which fail silently on the windows-latest -> windows-2025-vs2026
  rolling image.
- Add log-file-suffix to every install step. The action's docs note
  this is required for matrix builds; otherwise parallel jobs collide
  on the same log path, surfacing as "Path Validation Error: Path(s)
  specified in the action for caching do(es) not exist". This is
  exactly the failure annotation we hit.
- Pin Windows runner to windows-2022 (away from the rolling
  windows-latest label) and keep its install on the default `local`
  method - sub-package naming is undocumented for Windows so just
  install the full toolkit there. On Linux switch to `network`
  method with a minimal sub-package list (nvcc, cudart, curand,
  thrust) for ~5x faster installs.

Co-authored-by: Cursor <cursoragent@cursor.com>
The previous commit's switch to `method: network` with a custom
sub-packages list killed all three Linux jobs at install in ~14s -
at least one of the short package names (nvcc, cudart, cudart-dev,
curand, curand-dev, thrust) doesn't resolve to an apt package on
the NVIDIA CUDA repo for these toolkit versions.

Revert to the proven `local` + `--toolkit` pattern (skips samples
to keep install reasonable). Keep v0.2.35 of the action and the
log-file-suffix fix - those parts were correct.

Co-authored-by: Cursor <cursoragent@cursor.com>
CUDA 13.0 removed support for Maxwell, Pascal, and Volta
(sm_50..sm_72). The ci-multi-arch preset previously listed
70;80;90, which causes Configure to fail on Ubuntu 24.04 + CUDA
13.0.0 with `nvcc fatal: Unsupported gpu architecture 'compute_70'`.

Drop 70 from the CI preset; Ampere + Hopper is enough surface to
catch arch-specific compile bugs. Pre-Ampere coverage stays in
the `release-multi` preset (which devs use with CUDA <= 12).

Co-authored-by: Cursor <cursoragent@cursor.com>
The full local CUDA installer on windows-2022 silently hangs at
the install step (24+ minutes with no progress, well past the
known ~5-7 min upper bound). Switch to the action's `network`
method, which downloads ~500 MB of just the sub-packages we
link against:

  nvcc, cudart, curand, curand_dev, thrust,
  visual_studio_integration

Windows sub-package names use underscores (per the action's
windows-links.ts), distinct from Linux's hyphenated names.

Also add timeout-minutes (Linux 25, Windows 45) so any future
hang fails fast instead of holding the runner for 6 hours.

Drop sm_70 from the Windows configure step too, mirroring the
ci-multi-arch preset's Ampere+Hopper coverage (CUDA 13 is the
forcing function on Linux; Windows is on 12.6 here but staying
consistent).

Co-authored-by: Cursor <cursoragent@cursor.com>
* schedule: dedupe tile_traits/atom_traits across all algorithms

* schedule: introduce layout view, drop hardcoded CSR coupling

* schedule: parameterize setup by layout type, add ELL layout and SpMV examples

* docs: document layout contract and add custom_layout example

* spmv/ell_merge_path: drop redundant zero-init of y

* style: snake_case template params and Doxygen-friendly comments

* schedule: add flat_uniform_occupancy partitioner

* docs: move datasets, experimentation, reproducing-results into docs/

* docs: move detailed build instructions into docs/build.md

* build: swap GoogleTest fetch for Catch2 v3.7.1

* tests: layout-contract conformance for csr / ell / flat partitioner

* coo: fix csr_t conversion using wrong member name

* layout: add COO, CSC, BCSR, DIA formats and SpMV examples

* docs: tabulate the in-tree layouts in README
* headers: self-contain merge_path includes (thrust/device_vector, error, timer)

* tests: comprehensive Catch2 test suite (24 files, 81 cases)

* validation: SpMV-aware tolerance + deterministic seed + DIA size guard

* validation: rigorous f64-reference + Wilkinson per-row bound

* examples/spmv: build f32 + f64 variants per kernel

* ell/dia: size-probe guards on power-law and unstructured matrices

* cleanup: header self-contain + overflow-safe ell/dia size guards

Addresses review feedback on #36:

- reference.hxx, test_helpers.hxx: stop relying on transitive includes;
  pull in <cuda_runtime.h>, <algorithm>, <tuple> directly so the headers
  compile regardless of include order.
- test_util_math.cu: add <limits> for std::numeric_limits.
- ell_{thread_mapped,merge_path}.cu, dia_thread_mapped.cu: switch the
  preflight from "compute bytes, compare" to "compare cell count via
  division" so the multiplication can't overflow on hub-heavy / billion-
  row inputs. ELL now budgets sizeof(index_t) + sizeof(type_t) per cell
  (was sizeof(type_t) only, ~2x undercount); DIA also caps the
  diag_offsets allocation against the same 4 GiB limit.
- test_container_csr.cu: turn the previously-unused csr_pos vector into
  a per-slot claim flag so the COO->CSR round-trip test now also rejects
  duplicate emissions.

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
* market: mmap + from_chars reader, two-pass symmetric expansion

Rewrites the Matrix Market loader so loading com-Orkut.mtx (117 M lines,
234 M nnz after symmetric expansion) goes from ~5-8 minutes of fscanf
tokenization down to ~30-60 seconds, and the f64 path no longer trips
the sweep harness's 600 s timeout.

Loader path:
- Body parsing now operates on a memory-mapped view of the file via
  std::from_chars instead of fscanf. fscanf is locale-aware (a hot
  thread-local mutex on glibc) and serializes through stdio's userspace
  FILE buffer; std::from_chars is locale-independent and walks the
  mmap'd bytes directly. Measured ~30 M integers/sec vs. ~150 K/sec.
- Symmetric matrices use a two-pass strategy over the same mmap'd
  region: a structural pass counts off-diagonals (skipping the value
  column entirely), then a parsing pass emits straight into a correctly-
  sized COO with the off-diagonal mirror inlined. Page cache stays warm
  across the two passes, so wall cost is ~1.5x a single pass, not 2x.

Memory + correctness:
- Peak host RAM during a symmetric load drops to just the final COO.
  The previous loader allocated three full-sized intermediates, then
  copy-assigned them onto the final COO, doubling peak. On com-Orkut
  f64 that 1-2 GiB swing was the difference between "fits" and "swap".
- All counters are std::size_t. The previous int-typed off-diagonal
  counter would silently overflow on > 2^31 final-nnz matrices
  (Friendster-class), corrupting the symmetric expansion.

Drops the bundled NIST mmio library:
- include/loops/container/detail/mmio.{hxx,cpp} are removed; nothing
  else in the tree references them. Net -260 LOC of C-style FILE*
  parsing, replaced by ~400 LOC of header-only walk-pointer C++17.

New:
- detail/mapped_file.hxx -- portable RAII read-only mmap (POSIX +
  Windows MapViewOfFile), with MADV_SEQUENTIAL on the Linux side so
  cold-cache reads of multi-GiB files saturate readahead.
- detail/mtx_parser.hxx -- header-only parser primitives (skip_ws,
  parse_size_t, parse_double, parse_banner). std::from_chars where
  available, std::strtod fallback for older libc++ / MSVC.
- unittests/test_market_loader.cu -- coverage for general/symmetric/
  pattern/integer matrices, comment-line handling, and rejection of
  unsupported variants (complex/hermitian/skew/array/zero-indexed).

Public API is unchanged: matrix_market_t<index_t, offset_t, type_t>::
load(filename) still returns a host coo_t, .filename / .dataset are
still populated.

Co-authored-by: Cursor <cursoragent@cursor.com>

* test: fix nnz header in symmetric-expansion mtx fixture

The fixture declared "3 3 3" but supplied four body records, so the
loader (correctly) truncated at three and the assertions (incorrectly)
expected the fourth. Bump the header to "3 3 4" so the test exercises
all four records and the expected total of five expanded entries.

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Claude Opus 4 <noreply@anthropic.com>
Move build, datasets, experimentation, and reproducing-results docs
into site/content/ as the single source of truth. Delete docs/ dir,
DESIGN.md, and simplify build.py. README now links to the website.
…oad (#43)

* Fix CUDA warning: remove __device__ from container-based indices() overload

Agent-Logs-Url: https://github.com/gunrock/loops/sessions/f302ee1d-2ce1-45d6-aa26-5487cfec144e

Co-authored-by: neoblizz <9790745+neoblizz@users.noreply.github.com>

* range: fix no-op enable_if guard on indices() overload

std::enable_if<cond> without ::type names the trait struct, which is
well-formed regardless of cond, so the SFINAE guard never constrained
the overload; only the trailing return type incidentally did. Switch to
enable_if_t so has_size actually participates, and document why the
container overload is host-only.

Co-authored-by: Cursor <cursoragent@cursor.com>

* ci: cap build parallelism to stop runner OOM

The build matrix grew to ~29 Thrust/CUB-heavy targets across f32/f64 and
two SM arches. With unbounded make -j, every nvcc/cicc spawned at once and
exhausted the ~16 GB hosted runner; jobs died mid-compile with SIGTERM
(exit 143, "runner received a shutdown signal").

Cap the ci-multi-arch build preset to jobs: 2, drop the bare -j from the
Linux build step so the preset value applies, cap the Windows build at
-j 2, and raise the Linux job timeout 25 -> 40 to absorb the slower
serial-ish build.

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Co-authored-by: neoblizz <9790745+neoblizz@users.noreply.github.com>
Co-authored-by: neoblizz <osama94@gmail.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants