[CUB] Add runtime-sized shared-memory histogram privatization

[robobryce]

Why

CUB currently keeps at most 256 privatized histogram bins per channel in shared memory. Larger histograms use per-block global-memory privatization, even when the complete histogram would fit comfortably in the device's runtime-sized shared-memory budget.

The raw experiments in #10547 explored several ways to extend this range, including multiple compile-time bin tiers, global staging slabs, cooperative combine kernels, and hybrid paths. This PR extracts the smallest useful design from that work: one runtime-sized shared-memory kernel with a direct per-block merge into the output.

This PR is independent of the benchmark/testing PR #10548 and is based directly on main.

What changed

Runtime-sized privatized storage

AgentHistogram can now receive its privatized counter storage from an extern __shared__ allocation. The existing accumulation, zeroing, and output-store code continues to access the histogram through the same per-channel indexing interface.

The normal static-storage constructor is unchanged. The dynamic-storage constructor always selects the shared-memory path and initializes channel pointers into the contiguous runtime allocation.

One dynamic shared-memory kernel

The new sweep kernel keeps only the small agent metadata and block-load scratch storage in static shared memory. Histogram counters are allocated dynamically at launch according to the actual number of bins and sizeof(CounterT).

After processing its tiles, each block merges its private histogram directly into the output using the existing StoreOutput path. There is no global staging slab, follow-on combine kernel, cooperative launch, or cache algorithm in this PR.

Conservative dispatch

The new path is initially limited to single-channel, host-initialized, non-byte EVEN and RANGE histograms. The existing 256-bin static path remains unchanged. Histograms use the dynamic path only when:

• they have more than 256 bins; and
• their counters require no more than 32 KiB of dynamic shared memory per block.

The capacity check is byte-based, so a 64-bit counter supports half as many bins as a 32-bit counter. Larger or otherwise unsupported cases retain the existing global-memory privatized fallback.

The occupancy query receives the dynamic allocation size, and the temporary-storage calculation does not allocate an unused per-block global histogram slab for dynamic-shared-memory cells.

Deliberately excluded

This PR does not include:

• high-bin direct-atomic or caching algorithms;
• staging/combine or cooperative-grid implementations;
• multi-channel dynamic shared-memory selection;
• device-initialized C Parallel launcher support;
• architecture-specific tuning or low-bin RANGE policy changes;
• selector framework changes from the raw research branch.

Those concerns can be measured and reviewed independently after this basic shared-memory extension is accepted.

Tests

The histogram test adds single-channel EVEN and RANGE cases above the static 256-bin boundary. It covers:

• 1,024 and 4,096 bins;
• 32-bit counters, where both cases fit the dynamic budget; and
• 64-bit counters, including a case at the 32 KiB capacity boundary.

Formatting and source checks pass:

pre-commit run --files \
  cub/cub/agent/agent_histogram.cuh \
  cub/cub/device/dispatch/dispatch_histogram.cuh \
  cub/cub/device/dispatch/kernels/kernel_histogram.cuh \
  cub/test/catch2_test_device_histogram.cu

git diff --check upstream/main...HEAD

The focused test target builds successfully with CUDA 13.3.33 for SM90:

CUDAHOSTCXX=/usr/bin/c++ CXX=/usr/bin/c++ \
  cmake --fresh --preset cub-cpp20 -DCMAKE_CUDA_ARCHITECTURES=90

ninja -C build/cub-cpp20 cub.test.device.histogram

Runtime execution is blocked by the current machine's driver/toolkit combination. All three histogram launch modes, including unchanged tests, fail before assertions with:

cudaErrorUnsupportedPtxVersion: the provided PTX was compiled with an unsupported toolchain

Runtime validation therefore still needs a compatible CUDA driver for the CUDA 13.3.33-generated PTX.

Relationship to the research snapshot

This is a clean reimplementation of the final simplified shared-memory idea from #10547, rather than a chronological cherry-pick of the research commits. The earlier staging, cooperative, and fixed-tier implementations were intentionally discarded.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[robobryce]

Superseded by the same focused change under the corrected branch namespace: #10556.

[robobryce]

Replacement draft under the corrected branch namespace: #10556 ().

[robobryce]

Corrected branch name: pr/histocache/smem-privatized.