dylanyunlong/project_6
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity

feat(cccl): integrate missing CCCL directories — python/, ci/, .agent/, docs/, test/

Browse Source 
Sparse-checkout from NVIDIA/cccl main branch to complete cccl_upstream:

Added:
- python/cuda_cccl/ (226 files) — Python bindings for device-level algorithms
  Critical for muh toolchain: cuda.compute.reduce_into, scan, radix_sort, etc.
  Includes 204 .py files with full test coverage for all 27 algorithms
- ci/ (163 files) — Build/test infrastructure
  build_cub.sh, test_cub.sh, build_and_test_targets.sh, matrix.yaml
  Directly maps to our [INFRA-CI] and [INFRA-BUILD] items
- .agent/skills/ (7 files) — NVIDIA's own agent skills for CCCL
  cccl-style/SKILL.md, cccl-test/SKILL.md, sass-diff/SKILL.md
- docs/ (491 files) — Official CCCL documentation
  CI references, CMake guides, Python compute docs, libcudacxx PTX docs
- test/ (12 files) — Top-level integration tests (cuda_smoke, stdpar)
- Root configs: .clang-format, .clang-tidy, CONTRIBUTING.md, pyproject.toml
- CLAUDE.md symlink → AGENTS.md (NVIDIA's standard)

cccl_upstream now mirrors full NVIDIA/cccl structure:
  Before: 42M (cub + thrust + libcudacxx + cudax + c + examples + benchmarks)
  After:  53M (+python +ci +docs +.agent +test +configs)

This completes the CCCL base needed for:
- [muh-bench] items: ci/util/build_and_test_targets.sh for targeted builds
- [CCCL-verify] items: python/cuda_cccl/tests/ as reference implementations
- [CCCL-test] items: ci/test_cub.sh, ci/test_thrust.sh
- Agent workflow: .agent/skills/ for consistent style and test patterns
This commit is contained in:
muh-bot
2026-08-07 02:34:33 +00:00
parent 3f97dca7ad
commit 2a7ca101d7
908 changed files with 121615 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

183
cccl_upstream/docs/cub/Doxyfile Normal file
View File

@@ -0,0 +1,183 @@
# Doxyfile for CUB
PROJECT_NAME = CUB
OUTPUT_DIRECTORY = ../_build/doxygen/cub
CREATE_SUBDIRS = NO
GENERATE_HTML = NO
GENERATE_LATEX = NO
GENERATE_XML = YES
XML_OUTPUT = xml
XML_PROGRAMLISTING = YES
INPUT = ../../cub/cub \
../../cub/cub/thread \
../../cub/cub/warp \
../../cub/cub/block \
../../cub/cub/device \
../../cub/cub/device/dispatch/tuning \
../../cub/cub/grid \
../../cub/cub/iterator
RECURSIVE = YES
EXCLUDE_PATTERNS = */detail/* */dispatch/dispatch_* */dispatch/kernels/* */kernels/* */test/* */examples/*
EXCLUDE_SYMBOLS = *detail* CUB_DETAIL*
FILE_PATTERNS = *.cuh *.h
EXTENSION_MAPPING = cuh=C++ cu=C++
# Documentation extraction settings
EXTRACT_ALL = YES
EXTRACT_PRIVATE = NO
EXTRACT_STATIC = YES
EXTRACT_LOCAL_CLASSES = YES
HIDE_UNDOC_MEMBERS = NO
HIDE_UNDOC_CLASSES = NO
SHOW_INCLUDE_FILES = YES
INLINE_INHERITED_MEMB = YES
FULL_PATH_NAMES = YES
STRIP_FROM_PATH = ../../cub
STRIP_FROM_INC_PATH = ../../cub
SHORT_NAMES = NO
# Parsing settings
JAVADOC_AUTOBRIEF = YES
QT_AUTOBRIEF = NO
MULTILINE_CPP_IS_BRIEF = NO
INHERIT_DOCS = YES
SEPARATE_MEMBER_PAGES = NO
TAB_SIZE = 4
BUILTIN_STL_SUPPORT = YES
# Preprocessing
ENABLE_PREPROCESSING = YES
MACRO_EXPANSION = YES
EXPAND_ONLY_PREDEF = NO
SEARCH_INCLUDES = YES
SKIP_FUNCTION_MACROS = YES
# IMPORTANT: Aliases for custom commands
# The rst alias enables embedding reStructuredText in doxygen comments
ALIASES = "rst=\verbatim embed:rst:leading-asterisk"
ALIASES += "endrst=\endverbatim"
ALIASES += "rststar=\verbatim embed:rst:leading-asterisk"
ALIASES += "inlinerst=\verbatim embed:rst:inline"
# Key aliases that are used within @rst blocks (from repo.toml)
ALIASES += "smemwarpreuse=A subsequent ``__syncwarp()`` warp-wide barrier should be invoked after calling this method if the collective's temporary storage (e.g., ``temp_storage``) is to be reused or repurposed."
ALIASES += "smemreuse=A subsequent ``__syncthreads()`` threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., ``temp_storage``) is to be reused or repurposed."
ALIASES += "smemreuse{1}=After any operation, a subsequent ``__syncthreads()`` barrier is required if the collective's \1 is to be reused or repurposed"
ALIASES += "smemstorage{1}=The operations exposed by \1 require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the ``__shared__`` keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or ``union``'d with other storage allocation types to facilitate memory reuse."
ALIASES += "granularity=Efficiency is increased with increased granularity ``ITEMS_PER_THREAD``. Performance is also typically increased until the additional register pressure or shared memory allocation size causes SM occupancy to fall too low. Consider variants of ``cub::BlockLoad`` for efficiently gathering a :ref:`blocked arrangement <flexible-data-arrangement>` of elements across threads."
ALIASES += "blocksize=The number of threads in the block is a multiple of the architecture's warp size"
ALIASES += "ptxversion=The PTX compute capability for which to to specialize this collective, formatted as per the ``__CUDA_ARCH__`` macro (e.g., 750 for sm_75). Useful for determining the collective's storage requirements for a given device from the host. (Default: the value of ``__CUDA_ARCH__`` during the current compiler pass)"
ALIASES += "blockcollective{1}=Every thread in the block uses the \1 class by first specializing the \1 type, then instantiating an instance with parameters for communication, and finally invoking one or more collective member functions."
ALIASES += "warpcollective{1}=Every thread in the warp uses the \1 class by first specializing the \1 type, then instantiating an instance with parameters for communication, and finally invoking or more collective member functions."
ALIASES += "devicestorage=Temporary storage for this operation. If ``d_temp_storage`` is ``nullptr``, the required size is written to ``temp_storage_bytes`` without dereferencing iterators or launching kernels. Otherwise, ``d_temp_storage`` must point to a device-accessible allocation of at least ``temp_storage_bytes`` bytes. No special alignment is required. See :ref:`device-temp-storage` for usage guidance."
ALIASES += "devicestorageP=This operation requires a relatively small allocation of temporary device storage that is ``O(P)``, where ``P`` is the number of streaming multiprocessors on the device (and is typically a small constant relative to the input size ``N``)."
ALIASES += "devicestorageNP=This operation requires an allocation of temporary device storage that is ``O(N+P)``, where ``N`` is the length of the input and ``P`` is the number of streaming multiprocessors on the device."
ALIASES += "devicestorageNCP=This operation requires a relatively small allocation of temporary device storage that is ``O(N/C + P)``, where ``N`` is the length of the input, ``C`` is the number of concurrent threads that can be actively scheduled on each streaming multiprocessor (typically several thousand), and ``P`` is the number of streaming multiprocessors on the device."
ALIASES += "cdp_class{1}= - Dynamic parallelism. \1 methods can be called within kernel code on devices in which CUDA dynamic parallelism is supported."
ALIASES += "determinism{1}= - Determinism. The default reproducibility guarantee is ``\1``. A different guarantee can be requested through the execution environment with ``cuda::execution::require``. See :ref:`Determinism in CUB <cub-determinism>` for the supported guarantees."
ALIASES += "iterator=(may be a simple pointer type)"
ALIASES += "offset_size1=(Consider using 32-bit values as offsets/lengths/etc. For example, ``int`` will typically yield better performance than ``size_t`` in 64-bit memory mode.)"
ALIASES += "offset_size2=Careful consideration should be given to the size of integer types used for offsets and lengths. Many (if not most) scenarios will only require 32-bit offsets (e.g., ``int``). 64-bit offset types (e.g., ``size_t`` on 64-bit memory mode) can consume a significant amount of thread storage resources, adversely affecting processor occupancy and performance."
ALIASES += "rowmajor=For multi-dimensional blocks, threads are linearly ranked in row-major order."
ALIASES += "blocked=Assumes a :ref:`blocked arrangement <flexible-data-arrangement>` of (*block-threads* * *items-per-thread*) items across the thread block, where *thread*\ :sub:`i` owns the *i*\ :sup:`th` range of *items-per-thread* contiguous items. For multi-dimensional thread blocks, a row-major thread ordering is assumed."
ALIASES += "striped=Assumes a :ref:`striped arrangement <flexible-data-arrangement>` of (*block-threads* * *items-per-thread*) items across the thread block, where *thread*\ :sub:`i` owns items (*i*), (*i* + *block-threads*), ..., (*i* + (*block-threads* * (*items-per-thread* - 1))). For multi-dimensional thread blocks, a row-major thread ordering is assumed."
ALIASES += "warpstriped=Assumes a *warp-striped arrangement* of elements across threads, where warp\ :sub:`i` owns the *i*\ :sup:`th` range of (*warp-threads* * *items-per-thread*) contiguous items, and each thread owns items (*i*), (*i* + *warp-threads*), ..., (*i* + (*warp-threads* * (*items-per-thread* - 1)))."
ALIASES += "linear_performance{1}=The work-complexity of \1 as a function of input size is linear, resulting in performance throughput that plateaus with problem sizes large enough to saturate the GPU."
ALIASES += "plots_below=Performance plots for other scenarios can be found in the detailed method descriptions below."
ALIASES += "identityzero=This operation assumes the value of obtained by the ``T``'s default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value \"zero\" for addition."
ALIASES += "lookback=`decoupled look-back <https://research.nvidia.com/publication/single-pass-parallel-prefix-scan-decoupled-look-back>`_"
# Predefined macros (based on repo.toml doxygen_predefined)
PREDEFINED = __device__= \
__host__= \
__global__= \
__forceinline__= \
"__declspec(x)=" \
"__align__(x)=" \
__cccl_lib_mdspan \
"CUB_NAMESPACE_BEGIN=namespace cub {" \
"CUB_NAMESPACE_END=}" \
"CUB_NS_PREFIX=" \
"CUB_NS_POSTFIX=" \
"CUB_NS_QUALIFIER=cub::" \
"CUB_DETAIL_MAGIC_NS_BEGIN=" \
"CUB_DETAIL_MAGIC_NS_END=" \
"_CCCL_AND=&&" \
"_CCCL_CONCEPT=constexpr bool " \
"_CCCL_CONSTEXPR_FRIEND=friend " \
"_CCCL_CONSTEXPR_CXX20=constexpr" \
"_CCCL_CONSTEXPR_CXX23=constexpr" \
"_CCCL_CTK_AT_LEAST(x, y)=1" \
"_CCCL_CTK_BELOW(x, y)=0" \
"_CCCL_CUDACC_AT_LEAST(x, y)=1" \
"_CCCL_CUDACC_BELOW(x, y)=0" \
_CCCL_DEVICE= \
_CCCL_DIAG_PUSH= \
_CCCL_DIAG_POP= \
"_CCCL_DIAG_SUPPRESS_CLANG(x)=" \
"_CCCL_DIAG_SUPPRESS_GCC(x)=" \
"_CCCL_DIAG_SUPPRESS_MSVC(x)=" \
"_CCCL_DIAG_SUPPRESS_NVHPC(x)=" \
_CCCL_DOXYGEN_INVOKED \
_CCCL_EXEC_CHECK_DISABLE= \
_CCCL_FORCEINLINE= \
"_CCCL_GLOBAL_CONSTANT=inline constexpr" \
"_CCCL_HAS_CTK()=1" \
_CCCL_HIDE_FROM_ABI= \
_CCCL_HOST= \
_CCCL_HOST_DEVICE= \
"_CCCL_REQUIRES(x)= ::cuda::std::enable_if_t<x, int> = 0>" \
_CCCL_STD_VER=2020 \
_CCCL_SUPPRESS_DEPRECATED_PUSH= \
_CCCL_SUPPRESS_DEPRECATED_POP= \
"_CCCL_TEMPLATE(x)=template<x, " \
"_CCCL_TRAIT(x, y)=x<y>::value" \
"_CCCL_TRAILING_REQUIRES(x)=-> x requires " \
_CCCL_TYPE_VISIBILITY_DEFAULT= \
_CCCL_TYPE_VISIBILITY_HIDDEN= \
_CCCL_API=inline \
_CCCL_HOST_DEVICE_API=inline \
_CCCL_DEVICE_API=inline \
_CCCL_HOST_API=inline \
_CCCL_NODEBUG_API=inline \
_CCCL_NODEBUG_DEVICE_API=inline \
_CCCL_NODEBUG_HOST_API=inline \
_CCCL_TRIVIAL_API=inline \
_CCCL_TRIVIAL_DEVICE_API=inline \
_CCCL_TRIVIAL_HOST_API=inline \
_CCCL_VISIBILITY_DEFAULT= \
_CCCL_VISIBILITY_HIDDEN= \
_CCCL_LIFETIMEBOUND= \
_CCCL_TRY=try \
_CCCL_CATCH=catch \
"_CCCL_CATCH_ALL=catch (...)" \
"_CCCL_CATCH_FALLTHROUGH=" \
_CCCL_PUBLIC_API=inline \
_CCCL_PUBLIC_DEVICE_API=inline \
_CCCL_PUBLIC_HOST_API=inline \
"_CUDAX_CONSTEXPR_FRIEND=friend" \
"_LIBCUDACXX_HAS_SPACESHIP_OPERATOR()=1" \
"CCCL_DEPRECATED=" \
"CCCL_DEPRECATED_BECAUSE(x)=" \
"CCCL_IGNORE_DEPRECATED_CPP_DIALECT" \
"CUB_DISABLE_NAMESPACE_MAGIC" \
"CUB_IGNORE_NAMESPACE_MAGIC_ERROR" \
"CUB_RUNTIME_FUNCTION=" \
"THRUST_FWD(x)=x" \
"THRUST_NAMESPACE_BEGIN=namespace thrust {" \
"THRUST_NAMESPACE_END=}" \
"THRUST_PREVENT_MACRO_SUBSTITUTION" \
"_CCCL_HOSTED()=1" \
_CCCL_DOXYGEN_INVOKED
# Quiet mode
QUIET = YES
WARNINGS = YES
WARN_AS_ERROR = FAIL_ON_WARNINGS
WARN_IF_UNDOCUMENTED = NO
WARN_IF_DOC_ERROR = YES
WARN_IF_INCOMPLETE_DOC = NO
WARN_NO_PARAMDOC = NO

376
cccl_upstream/docs/cub/benchmarking.rst Normal file
View File

@@ -0,0 +1,376 @@
Benchmarks
*************************************
.. TODO(bgruber): this guide applies to Thrust as well. We should rename it to "CCCL Benchmarks" and move it out of CUB
CUB comes with a set of `NVBench <https://github.com/NVIDIA/nvbench>`_-based benchmarks for its algorithms,
which can be used to measure the performance of CUB on your system on a variety of workloads.
The integration with NVBench allows to archive and compare benchmark results,
which is useful for continuous performance testing, detecting regressions, tuning, and optimization.
This guide gives an introduction into CUB's benchmarking infrastructure.
Building benchmarks
--------------------------------------------------------------------------------
CUB benchmarks are build as part of the CCCL CMake infrastructure.
Starting from scratch:
.. code-block:: bash
git clone https://github.com/NVIDIA/cccl.git
cd cccl
mkdir build
cd build
cmake .. --preset=benchmark
You clone the repository, create a build directory and configure the build with CMake.
The preset `benchmark` takes care of everything.
.. TODO(bgruber): do we have a public NVIDIA maintained table I can link here instead?
We use Ninja as CMake generator in this guide, but you can use any other generator you prefer.
You can then proceed to build the benchmarks.
You can list the available cmake build targets with, if you intend to only build selected benchmarks:
.. code-block:: bash
ninja -t targets | grep '\.bench\.'
cub.bench.adjacent_difference.subtract_left.base: phony
cub.bench.copy.memcpy.base: phony
...
cub.bench.transform.babelstream3.base: phony
cub.bench.transform_reduce.sum.base: phony
We also provide a target to build all benchmarks:
.. code-block:: bash
ninja cub.all.benches
.. _cub-benchmarking-running:
Running a benchmark
--------------------------------------------------------------------------------
After we built a benchmark, we can run it as follows:
.. code-block:: bash
./bin/cub.bench.adjacent_difference.subtract_left.base\
-d 0\
--stopping-criterion entropy\
--json base.json\
--md base.md
In this command, `-d 0` indicates that we want to run on GPU 0 on our system.
Setting `--stopping-criterion entropy` is advisable since it reduces runtime
and increase confidence in the resulting data.
It's not set as default yet, because NVBench is still evaluating it.
By default, NVBench will print the benchmark results to the terminal as Markdown.
`--json base.json` will save the detailed results in a JSON file as well for later use.
`--md base.md` will save the Markdown output to a file as well,
so you can easily view the results later without having to parse the JSON.
More information on what command line options are available can be found in the
`NVBench documentation <https://github.com/NVIDIA/nvbench/blob/main/docs/cli_help.md>`__.
The expected terminal output is something along the following lines (also saved to `base.md`),
shortened for brevity:
.. code-block:: bash
# Log
Run: [1/8] base [Device=0 T{ct}=I32 OffsetT{ct}=I32 Elements{io}=2^16]
Pass: Cold: 0.004571ms GPU, 0.009322ms CPU, 0.00s total GPU, 0.01s total wall, 334x
Run: [2/8] base [Device=0 T{ct}=I32 OffsetT{ct}=I32 Elements{io}=2^20]
Pass: Cold: 0.015161ms GPU, 0.023367ms CPU, 0.01s total GPU, 0.02s total wall, 430x
...
# Benchmark Results
| T{ct} | OffsetT{ct} | Elements{io} | Samples | CPU Time | Noise | GPU Time | Noise | Elem/s | GlobalMem BW | BWUtil |
|-------|-------------|------------------|---------|------------|---------|------------|--------|---------|--------------|--------|
| I32 | I32 | 2^16 = 65536 | 334x | 9.322 us | 104.44% | 4.571 us | 10.87% | 14.337G | 114.696 GB/s | 14.93% |
| I32 | I32 | 2^20 = 1048576 | 430x | 23.367 us | 327.68% | 15.161 us | 3.47% | 69.161G | 553.285 GB/s | 72.03% |
...
If you are only interested in a subset of workloads, you can restrict benchmarking as follows:
.. code-block:: bash
./bin/cub.bench.adjacent_difference.subtract_left.base ...\
-a 'T{ct}=I32'\
-a 'OffsetT{ct}=I32'\
-a 'Elements{io}[pow2]=[24,28]'\
The `-a` option allows you to restrict the values for each axis available for the benchmark.
See the `NVBench documentation <https://github.com/NVIDIA/nvbench/blob/main/docs/cli_help_axis.md>`__.
for more information on how to specify the axis values.
If the specified axis does not exist, the benchmark will terminate with an error.
If you want to plot the benchmark results, you can use the following script:
.. code-block:: bash
PYTHONPATH=./_deps/nvbench-src/python/scripts ./_deps/nvbench-src/python/scripts/nvbench_plot_bwutil.py base.json
The `-a` option is supported to restrict the values for some axes as well,
which is useful if you want to plot only a subset of workloads.
Use the `-b` option to select a specific benchmark by name
in case your JSON file contains results for multiple benchmarks.
Multiple benchmarks are selected by repeating the `-b` option.
.. code-block:: bash
PYTHONPATH=./_deps/nvbench-src/python/scripts ./_deps/nvbench-src/python/scripts/nvbench_plot_bwutil.py \
-b base -a Elements{io}[pow2]=28 base.json
.. _cub-benchmarking-comparing:
Comparing benchmark results
--------------------------------------------------------------------------------
Let's say you have a modification that you'd like to benchmark.
To compare the performance you have to build and run the benchmark as described above for the unmodified code,
saving the results to a JSON file, e.g. `base.json`.
Then, you apply your code changes (e.g., switch to a different branch, git stash pop, apply a patch file, etc.),
rebuild and rerun the benchmark, saving the results to a different JSON file, e.g. `new.json`.
You can now compare the two result JSON files using, assuming you are still in your build directory:
.. code-block:: bash
PYTHONPATH=./_deps/nvbench-src/python/scripts ./_deps/nvbench-src/python/scripts/nvbench_compare.py base.json new.json
The `PYTHONPATH` environment variable may not be necessary in all cases.
The script will print a Markdown report showing the runtime differences between each variant of the two benchmark run.
This could look like this, again shortened for brevity:
.. code-block:: bash
| T{ct} | OffsetT{ct} | Elements{io} | Ref Time | Ref Noise | Cmp Time | Cmp Noise | Diff | %Diff | Status |
|---------|---------------|----------------|------------|-------------|------------|-------------|------------|---------|----------|
| I32 | I32 | 2^16 | 4.571 us | 10.87% | 4.096 us | 0.00% | -0.475 us | -10.39% | FAIL |
| I32 | I32 | 2^20 | 15.161 us | 3.47% | 15.143 us | 3.55% | -0.018 us | -0.12% | PASS |
...
In addition to showing the absolute and relative runtime difference,
NVBench reports the noise of the measurements,
which corresponds to the relative standard deviation.
It then reports with statistical significance in the `Status` column
how the runtime changed from the base to the new version.
You can reduce the output to runs with larger differences using the `--threshold-diff` option,
passing the minimum percentage for a run to be shown, e.g., `0.05` for 5%.
.. code-block:: bash
PYTHONPATH=./_deps/nvbench-src/python/scripts ./_deps/nvbench-src/python/scripts/nvbench_compare.py \
--threshold-diff 0.05 base.json new.json
You can also plot the comparison by adding the `--plot` argument.
It's reasonable to combine this with the `-a` option again
to restrict the values for some axes.
.. code-block:: bash
PYTHONPATH=./_deps/nvbench-src/python/scripts ./_deps/nvbench-src/python/scripts/nvbench_compare.py \
-a Elements{io}[pow2]=28 --plot base.json new.json
Running all benchmarks directly from the command line
--------------------------------------------------------------------------------
To get a full snapshot of CUB's performance, you can run all benchmarks and save the results.
For example, inside a build directory you can run:
.. code-block:: bash
ninja cub.all.benches
benchmarks=$(ls bin | grep cub.bench); n=$(echo $benchmarks | wc -w); i=1; \
for b in $benchmarks; do \
echo "=== Running $b ($i/$n) ==="; \
./bin/$b -d 0 --stopping-criterion entropy --json $b.json --md $b.md; \
((i++)); \
done
This will generate one JSON and one Markdown file for each benchmark.
You can archive those files for later comparison or analysis.
Running all benchmarks via tuning scripts (alternative)
--------------------------------------------------------------------------------
The benchmark suite can also be run using the :ref:`tuning infrastructure <cub-tuning-infra>`.
The tuning infrastructure handles building benchmarks itself, because it records the build times.
Therefore, it's critical that you run it in a clean build directory without any build artifacts.
Running cmake is enough. Alternatively, you can also clean your build directory.
Furthermore, the tuning scripts require some additional python dependencies, which you have to install:
.. code-block:: bash
ninja clean
pip install --user fpzip pandas scipy
To select the appropriate CUDA GPU, first identify the GPU ID by running `nvidia-smi`, then set the
desired GPU using `export CUDA_VISIBLE_DEVICES=x <https://docs.nvidia.com/cuda/cuda-c-programming-guide/#cuda-environment-variables>`_,
where `x` is the ID of the GPU you want to use (e.g., `1`).
This ensures your application uses only the specified GPU.
We can then run the full benchmark suite from the build directory with:
.. code-block:: bash
export CUDA_VISIBLE_DEVICES=0 # or any other GPU ID
PYTHONPATH=../benchmarks/scripts ../benchmarks/scripts/run.py
You can expect the output to look like this:
.. code-block:: bash
&&&& RUNNING bench
ctk: 12.2.140
cub: 812ba98d1
&&&& PERF cub_bench_adjacent_difference_subtract_left_base_T_ct__I32___OffsetT_ct__I32___Elements_io__pow2__16 4.095999884157209e-06 -sec
&&&& PERF cub_bench_adjacent_difference_subtract_left_base_T_ct__I32___OffsetT_ct__I32___Elements_io__pow2__20 1.2288000107218977e-05 -sec
&&&& PERF cub_bench_adjacent_difference_subtract_left_base_T_ct__I32___OffsetT_ct__I32___Elements_io__pow2__24 0.00016998399223666638 -sec
&&&& PERF cub_bench_adjacent_difference_subtract_left_base_T_ct__I32___OffsetT_ct__I32___Elements_io__pow2__28 0.002673664130270481 -sec
...
The tuning infrastructure will build and execute all benchmarks and their variants one after each other,
reporting the time in seconds it took to execute the benchmarked region.
It's also possible to benchmark a subset of algorithms and workloads, by running in a build directory:
.. code-block:: bash
export CUDA_VISIBLE_DEVICES=0 # or any other GPU ID
PYTHONPATH=../benchmarks/scripts ../benchmarks/scripts/run.py -R '.*scan.exclusive.sum.*' -a 'Elements{io}[pow2]=[24,28]' -a 'T{ct}=I32'
&&&& RUNNING bench
ctk: 12.6.77
cccl: v2.7.0-rc0-265-g32aa6aa5a
&&&& PERF cub_bench_scan_exclusive_sum_base_T_ct__I32___OffsetT_ct__U32___Elements_io__pow2__28 0.003194367978721857 -sec
&&&& PERF cub_bench_scan_exclusive_sum_base_T_ct__I32___OffsetT_ct__U64___Elements_io__pow2__28 0.00319383991882205 -sec
&&&& PASSED bench
The `-R` option allows you to specify a regular expression for selecting benchmarks.
The `-a` restricts the values for an axis across all benchmarks
See the `NVBench documentation <https://github.com/NVIDIA/nvbench/blob/main/docs/cli_help_axis.md>`__.
for more information on how to specify the axis values.
Contrary to running a benchmark directly,
the tuning infrastructure will just ignore an axis value if a benchmark does not support,
run the benchmark regardless, and continue.
The tuning infrastructure stores results in an SQLite database called :code:`cccl_meta_bench.db` in the build directory.
This database persists across tuning runs.
If you interrupt the benchmark script and then launch it again, only missing benchmark variants will be run.
Comparing results of multiple tuning databases
--------------------------------------------------------------------------------
Benchmark results captured in different tuning databases can be compared as well:
.. code-block:: bash
<cccl_git_root>/benchmarks/scripts/compare.py -o cccl_meta_bench1.db cccl_meta_bench2.db
This will print a Markdown report showing the runtime differences and noise for each variant.
Furthermore, you can plot the results, which requires additional python packages:
.. code-block:: bash
pip install fpzip pandas matplotlib seaborn tabulate PyQt5 colorama
You can plot one or more tuning databases as a bar chart or a box plot (add `--box`):
.. code-block:: bash
<cccl_git_root>/benchmarks/scripts/sol.py cccl_meta_bench.db ...
This is useful to display the current performance of CUB as captured in a single tuning database,
or visually compare the performance of CUB across different tuning databases
(from different points in time, on different GPUs, etc.).
Dumping benchmark results from a tuning database
--------------------------------------------------------------------------------
The resulting database contains all samples, which can be extracted into JSON files:
.. code-block:: bash
<cccl_git_root>/benchmarks/scripts/analyze.py -o ./cccl_meta_bench.db
This will create a JSON file for each benchmark variant next to the database.
For example:
.. code-block:: bash
cat cub_bench_scan_exclusive_sum_base_T_ct__I32___OffsetT_ct__U32___Elements_io__pow2__28.json
[
{
"variant": "base ()",
"elapsed": 2.6299014091,
"center": 0.003194368,
"bw": 0.8754671386,
"samples": [
0.003152896,
0.0031549439,
...
],
"Elements{io}[pow2]": "28",
"base_samples": [
0.003152896,
0.0031549439,
...
],
"speedup": 1
}
]
Profiling benchmarks with Nsight Compute
--------------------------------------------------------------------------------
If you want to see profiling metrics on source code level,
you have to recompile your benchmarks with the `-lineinfo` option.
With cmake, you can just add `-DCMAKE_CUDA_FLAGS=-lineinfo` when invoking cmake in the `build` directory:
.. code-block:: bash
cmake .. --preset=benchmark -DCMAKE_CUDA_FLAGS=-lineinfo
To profile the kernels, use the `ncu` command.
A typical invocation, if you work on a remote cluster, could look like this:
.. code-block:: bash
ncu --set full --import-source yes -o base.ncu-rep -f ./bin/thrust.bench.transform.basic.base -d 0 --profile
The option `--set full` instructs `ncu` to collect all metrics.
This requires rerunning some kernels and takes more time.
`--import-source yes` imports the source code into the report file,
so you can see metrics not only in SASS but also in your source code,
even if you copy the resulting report away from the source code.
`-o base.ncu-rep` specifies the output file and `-f` overwrites the output file if it already exists.
`--profile` tells NVBench to run only one iteration, which speeds up profiling.
For inspecting the profiling report, we recommend using the GUI of Nsight Compute.
If you run on a remote machine, you may want to copy the report `base.ncu-rep` back to your local workstation,
before viewing the report using `ncu-ui`:
.. code-block:: bash
scp <remote hostname>:<cccl repo directory>/build/base.ncu-rep .
ncu-ui base.ncu-rep
The version of `ncu-ui` needs to be at least as high as the version of `ncu` used to create the report.
Authoring benchmarks
--------------------------------------------------------------------------------
CUB's benchmarks serve a dual purpose.
They are used to measure and compare the performance of CUB and to tune CUB's algorithms.
More information on how to create new benchmarks is provided in the :ref:`CUB tuning infrastructure guide <cub-tuning-infra>`.

26
cccl_upstream/docs/cub/block_wide.rst Normal file
View File

@@ -0,0 +1,26 @@
.. _block-module:
Block-Wide "Collective" Primitives
==================================================
.. toctree::
:glob:
:hidden:
:maxdepth: 2
api/block
CUB block-level algorithms are specialized for execution by threads in the same CUDA thread block:
* :cpp:class:`cub::BlockAdjacentDifference` computes the difference between adjacent items partitioned across a CUDA thread block
* :cpp:class:`cub::BlockDiscontinuity` flags discontinuities within an ordered set of items partitioned across a CUDA thread block
* :cpp:struct:`cub::BlockExchange` rearranges data partitioned across a CUDA thread block
* :cpp:class:`cub::BlockHistogram` constructs block-wide histograms from data samples partitioned across a CUDA thread block
* :cpp:class:`cub::BlockLoad` loads a linear segment of items from memory into a CUDA thread block
* :cpp:class:`cub::BlockMergeSort` sorts items partitioned across a CUDA thread block
* :cpp:class:`cub::BlockRadixSort` sorts items partitioned across a CUDA thread block using radix sorting method
* :cpp:struct:`cub::BlockReduce` computes reduction of items partitioned across a CUDA thread block
* :cpp:class:`cub::BlockRunLengthDecode` decodes a run-length encoded sequence partitioned across a CUDA thread block
* :cpp:struct:`cub::BlockScan` computes a prefix scan of items partitioned across a CUDA thread block
* :cpp:struct:`cub::BlockShuffle` shifts items partitioned across a CUDA thread block
* :cpp:class:`cub::BlockStore` stores items partitioned across a CUDA thread block to a linear segment of memory

76
cccl_upstream/docs/cub/determinism.rst Normal file
View File

@@ -0,0 +1,76 @@
.. _cub-determinism:
Determinism
===============
Several ``cub`` device algorithms let you request a reproducibility guarantee for a call. The concepts
behind the three guarantees — ``not_guaranteed``, ``run_to_run``, and ``gpu_to_gpu`` — and the meaning
of *reproducibility* are described in the :ref:`CCCL determinism overview <cccl-determinism>`. This
page documents how to request a guarantee for a CUB algorithm and which algorithms support which
guarantees.
Requesting a guarantee
----------------------
A determinism guarantee is passed to a device algorithm through its execution environment using
``cuda::execution::require``. The example below requests run-to-run reproducibility for
``cub::DeviceReduce::Sum``:
.. literalinclude:: ../../cub/test/catch2_test_device_reduce_env_api.cu
:language: c++
:dedent:
:start-after: example-begin sum-env-determinism
:end-before: example-end sum-env-determinism
The general rules for requesting a guarantee are described in the
:ref:`CCCL determinism overview <cccl-determinism>`.
Each CUB algorithm has its own default guarantee, applied when none is requested, and its own type and
operator constraints for each guarantee, summarized below. Requesting a guarantee that an algorithm
does not support is rejected at compile time.
Support matrix
--------------
.. list-table::
:header-rows: 1
:widths: 30 18 18 18 16
* - Algorithm
- ``not_guaranteed``
- ``run_to_run``
- ``gpu_to_gpu``
- Default
* - ``cub::DeviceReduce`` (``Reduce``, ``Sum``, ``Min``, ``Max``, ``TransformReduce``, ...)
- Yes
- Yes
- Yes (partial)
- ``run_to_run``
* - ``cub::DeviceScan`` (``ExclusiveSum``, ``ExclusiveScan``, ``InclusiveSum``, ``InclusiveScan``, ...)
- Yes
- Yes (partial)
- Yes (partial)
- ``not_guaranteed``
* - ``cub::DeviceSegmentedReduce``
- Yes
- Yes
- No
- ``run_to_run``
.. note::
The set of algorithms that accept determinism requirements, and the type/operator constraints for
each guarantee, are expanding over time. The matrix above reflects the current implementation.
Algorithm-specific determinism models
--------------------------------------
The three guarantees describe the *scope* of reproducibility and fit most algorithms, where a
reproducible result means a *bitwise-identical* output. A few algorithms still use the same three
levels but extend the model with additional, algorithm-specific controls, documented on their own
pages:
- :ref:`cub::DeviceTopK <cub-topk-requirements>` — determinism applies to *set membership* (which
*K* items are selected) rather than a bitwise-identical buffer, and it adds tie-breaking
(``cuda::execution::tie_break``) and output-ordering (``cuda::execution::output_ordering``)
controls.

79
cccl_upstream/docs/cub/developer/block_scope.rst Normal file
View File

@@ -0,0 +1,79 @@
.. _cub-developer-guide-block-scope:
Block-scope
************
Overview
=========
Block-scope algorithms are provided by structures as well:
.. code-block:: c++
template <typename T,
int BLOCK_DIM_X,
BlockReduceAlgorithm ALGORITHM = BLOCK_REDUCE_WARP_REDUCTIONS,
int BLOCK_DIM_Y = 1,
int BLOCK_DIM_Z = 1>
class BlockReduce {
public:
struct TempStorage : Uninitialized<_TempStorage> {};
// (1) new constructor
__device__ __forceinline__ BlockReduce()
: temp_storage(PrivateStorage()),
linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z)) {}
__device__ __forceinline__ BlockReduce(TempStorage &temp_storage)
: temp_storage(temp_storage.Alias()),
linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z)) {}
};
While warp-scope algorithms only provide a single constructor that requires the user to provide temporary storage,
block-scope algorithms provide two constructors:
#. The default constructor that allocates the required shared memory internally.
#. The constructor that requires the user to provide temporary storage as argument.
In the case of the default constructor,
the block-level algorithm uses the ``PrivateStorage()`` member function to allocate the required shared memory.
This ensures that shared memory required by the algorithm is only allocated when the default constructor is actually called in user code.
If the default constructor is never called,
then the algorithm will not allocate superfluous shared memory.
.. code-block:: c++
__device__ __forceinline__ _TempStorage& PrivateStorage()
{
__shared__ _TempStorage private_storage;
return private_storage;
}
The ``__shared__`` memory has static semantic, so it's safe to return a reference here.
Specialization
====================================
Block-scope facilities usually expose algorithm selection to the user.
The algorithm is represented by the enumeration part of the API.
For the reduction case,
``BlockReduceAlgorithm`` is provided.
Specializations are stored in the ``cub/block/specializations`` directory.
Temporary storage usage
====================================
For block-scope algorithms,
it's unsafe to use temporary storage without synchronization:
.. code-block:: c++
using BlockReduce = cub::BlockReduce<int, 128> ;
__shared__ BlockReduce::TempStorage temp_storage;
int aggregate_1 = BlockReduce(temp_storage).Sum(thread_data_1);
// illegal, has to add `__syncthreads` between the two
int aggregate_2 = BlockReduce(temp_storage).Sum(thread_data_2);
// illegal, has to add `__syncthreads` between the two
foo(temp_storage);

530
cccl_upstream/docs/cub/developer/device_scope.rst Normal file
View File

@@ -0,0 +1,530 @@
.. _cub-developer-guide-device-scope:
Device-scope
*************
Overview
=========
Device-scope functionality is provided by classes called ``DeviceAlgorithm``,
where ``Algorithm`` is the implemented algorithm.
These classes then contain static member functions providing corresponding API entry points.
For example, device-level reduce will look like `cub::DeviceReduce::Sum`.
Here is a generic example:
.. code-block:: c++
struct DeviceAlgorithm {
// two step API
template <typename ...>
static cudaError_t Algorithm(void *d_temp_storage, size_t &temp_storage_bytes, ..., cudaStream_t stream = 0) {
// optional: minimal argument checking or setup to call dispatch layer
return detail::algorithm::dispatch(d_temp_storage, temp_storage_bytes, ..., stream);
}
// environment API
template <typename ..., typename Env = cuda::std::execution::env<>>
static cudaError_t Algorithm(..., const Env& env = {}) {
// optional: minimal argument checking or setup to call dispatch layer
using default_policy_selector = detail::algorithm::policy_selector_from_types<...>;
return dispatch_with_env_and_tuning<default_policy_selector>(
env, [&](auto policy_selector, void* storage, size_t& bytes, auto stream) {
return detail::algorithm::dispatch(d_temp_storage, temp_storage_bytes, ..., stream, policy_selector);
});
}
};
Device-scope APIs come in two flavors, two step APIs and environment APIs,
both return ``cudaError_t`` and take algorithm specific arguments.
The two step API accepts a ``stream`` as the last parameter (``NULL`` stream by default)
and the first two parameters are always ``void *d_temp_storage, size_t &temp_storage_bytes``.
The environment API just takes an environment as the last parameter (empty environment by default).
The implementation may consist of some minimal argument checking, but should forward as soon as possible to the dispatch layer.
Device-scope algorithms are implemented in files located in `cub/device/device_***.cuh`.
The two step API is called in two phases:
1. Temporary storage size is calculated and returned in ``size_t &temp_storage_bytes``.
2. ``temp_storage_bytes`` of memory is expected to be allocated and ``d_temp_storage`` is expected to be the pointer to this memory.
The following example illustrates this pattern:
.. code-block:: c++
// First call: Determine temporary device storage requirements
std::size_t temp_storage_bytes = 0;
cub::DeviceReduce::Sum(/*d_temp_storage*/ nullptr, temp_storage_bytes, d_in, d_out, num_items);
// Allocate temporary storage
thrust::device_vector<unsigned char> temp_storage(temp_storage_bytes, thrust::no_init);
// Second call: Perform algorithm
cub::DeviceReduce::Sum(temp_storage.data().get(), temp_storage_bytes, d_in, d_out, num_items);
.. warning::
Even if the algorithm doesn't need temporary storage as scratch space,
the overload with ``void *d_temp_storage, size_t &temp_storage_bytes``
still requires one byte of memory to be allocated.
The environment overloads just require a single call, but setting up the environment may be more complex:
.. code-block:: c++
// Setup environment, everything is optional
auto env = cuda::std::execution::env{
stream_ref,
memory_resource,
cuda::execution::require(requirements...),
cuda::execution::tune(policy_selectors....)
};
// Perform algorithm
cub::DeviceReduce::Sum(d_in, d_out, num_items, env);
The device layer handles the extraction of the:
* CUDA stream
* memory resource
* requirements
* guarantees (TODO(bgruber): are those public?)
* :ref:`tuning policy selectors <cub-policy-selectors>`
from the environment argument, or provide default values in case the environment does not contain them.
They typically use helper functions like ``dispatch_with_env`` and ``dispatch_with_env_and_tuning``.
Some CUB APIs require no temporary storage and may omit the ``void *d_temp_storage, size_t &temp_storage_bytes`` parameters.
Their environment overloads will also ignore any passed memory resource.
Dispatch layer
====================================
A dispatch function exists for each device-scope algorithm (e.g., ``detail::reduce::dispatch``),
and is located in ``cub/device/dispatch``.
Most device-scope algorithms share one dispatch function.
Only device-scope algorithms have dispatch functions, which are also referred to as the dispatch layer.
The dispatch layer follows a certain architecture.
The high-level control flow is represented by the code below.
A more precise description is given later.
..
TODO(bgruber): consider removing the numbers below. control flow is now easier to follow
.. code-block:: c++
// Device-scope API
cudaError_t cub::DeviceAlgorithm::Algorithm(d_temp_storage, temp_storage_bytes, ...) {
return detail::algorithm::dispatch(d_temp_storage, temp_storage_bytes, ...);
}
namespace detail::algorithm {
cudaError_t dispatch(
void *d_temp_storage, size_t &temp_storage_bytes,
...,
cudaStream_t stream, PolicySelector policy_selector = {}) {
cuda::compute_capability cc{};
ptx_compute_cap(cc);
const /*or constexpr*/ AlgorithmPolicy active_policy = policy_selector(cc);
// host-side implementation of algorithm, calls kernels
kernel<PolicySelector><<<grid_size, active_policy.threads_per_block>>>(...);
}
template <typename PolicySelector>
__launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
void kernel(...) {
static constexpr auto policy = current_policy<PolicySelector>();
using agent_policy = AgentPolicy<policy.threads_per_block, policy.items_per_thread, ...>;
using agent = AgentAlgorithm<agent_policy, InputIteratorT, OffsetT, ...>;
agent a{...};
a.Process();
}
template <int ThreadsPerBlock, ...>
struct AgentPolicy { // legacy
static constexpr threads_per_block = ThreadsPerBlock;
...
};
template <typename Policy, ...>
struct AlgorithmAgent {
void Process() { ... }
};
}
Let's look at each of the building blocks closer.
The dispatch function
------------------------------------
The dispatch function is typically a simple function template called ``dispatch`` inside an algorithm-specific namespace.
It basically receives the same parameters as the public API entry point,
but they may have already been modified, extended, or generalized (to map many public APIs to the same dispatch function).
The goal of the dispatch function is to setup the execution of the algorithm on the device
by selecting the appropriate policy for the current GPU's compute capability,
preparing temporary storage and shared memory, configuring kernel launches, launching kernels, etc.
There are two style of dispatch functions, depending on whether the policy is needed as runtime or compile-time value:
.. code-block:: c++
namespace detail::algorithm {
// Dispatch version A - runtime policy
cudaError_t dispatch(
void *d_temp_storage, size_t &temp_storage_bytes,
...,
cudaStream_t stream, PolicySelector policy_selector = {}) {
cuda::compute_capability cc{};
ptx_compute_cap(cc);
const auto active_policy = policy_selector(cc); // runtime-time policy
// host-side implementation of algorithm, calls kernels
kernel<PolicySelector><<<grid_size, active_policy.threads_per_block>>>(...);
}
}
The dispatch function starts by querying the target compute capability for which compiled GPU code (PTX or SASS) is available,
by calling ``ptx_compute_cap``.
If the host code does not require the policy for this compute capability at compile-time (version A),
we can just pass the compute capability to the :ref:`policy selector <cub-policy-selectors>` to obtain the tuning policy at runtime.
The values from the policy are then used to setup resources and the kernel.
The kernel is then instantiated using only the type of the policy selector (not a concrete tuning policy),
so there is only one kernel instantiation across all target architectures compiled for.
More on that later.
If the host code needs the policy as a compile-time value (version B), we have to use ``dispatch_compute_cap``.
dispatch_compute_cap
------------------------------------
``dispatch_compute_cap`` maps a runtime ``cuda::compute_capability`` to a compile-time policy value
and calls the user-provided functor ``f`` with a nullary callable (a ``policy_getter``)
that returns the policy as a compile-time constant.
The policy getter is necessary to work around a C++17 limitation.
.. code-block:: c++
namespace detail::algorithm {
// Dispatch version B - compile-time policy
cudaError_t dispatch(
void *d_temp_storage, size_t &temp_storage_bytes,
...,
cudaStream_t stream, PolicySelector policy_selector = {}) {
cuda::compute_capability cc{};
ptx_compute_cap(cc);
return dispatch_compute_cap(policy_selector, cc, [&](auto policy_getter) {
constexpr auto active_policy = policy_getter(); // compile-time policy
static_assert(active_policy.tile_size() * sizeof(T) <= 48 * 1024, "Not enough SMEM");
// host-side implementation of algorithm, calls kernels
kernel<PolicySelector><<<grid_size, active_policy.threads_per_block>>>(...);
});
}
template <typename PolicySelector, typename F>
cudaError_t dispatch_compute_cap(PolicySelector, cuda::compute_capability cc, F&& f) { // (2)
// fold over __CUDA_ARCH_LIST__, calling f with a policy_getter
// that returns the policy for the matching arch as a compile-time value
}
}
Inside the lambda, ``policy_getter()`` returns the selected policy as a constant expression,
so compile-time branching (i.e. ``if constexpr``) or static assertions using policy values is possible.
Internally, ``dispatch_compute_cap`` uses ``__CUDA_ARCH_LIST__`` / ``NV_TARGET_SM_INTEGER_LIST`` (or all known compute capabilities as fallback)
to create one instantiation of ``f`` per distinct value of ``policy_selector(cc)`` (not per compute capability).
This results in one template instantiation of ``f`` per distinct policy value.
The kernel is then again only instantiated once using the type of the policy selector.
Kernels
------------------------------------
Kernels are templated on the ``PolicySelector`` type,
which is stateless and the same for all target architectures compiled for.
.. code-block:: c++
namespace detail::algorithm {
template <typename PolicySelector, typename InputIteratorT, typename OffsetT, /* ... */>
__launch_bounds__(int(current_policy<PolicySelector>().reduce.threads_per_block))
void DeviceReduceKernel(InputIteratorT d_in, OffsetT num_items, /* ... */) {
static constexpr auto policy = current_policy<PolicySelector>();
using agent_policy = AgentPolicy<policy.threads_per_block, policy.items_per_thread, ...>;
using agent = AgentAlgorithm<agent_policy, InputIteratorT, OffsetT, ...>;
agent a{...};
a.Process();
}
template <typename PolicySelector>
constexpr auto current_policy() {
return PolicySelector{}(cuda::compute_capability{__CUDA_ARCH__ / 10}); // simplified
}
}
``PolicySelector`` must be stateless (``is_empty_v<PolicySelector>`` is ``true``),
so it can be default-constructed in device code where needed.
The utility function ``current_policy`` can only be called in device code.
It selects the target compute capability based on compiler macros of the current device compilation pass
and retrieves a tuning policy from the policy selector.
The kernel typically uses ``current_policy`` in two places,
to get the block size to define the launch bounds,
and inside the kernel to setup various sub algorithms (like agents).
Because C++17 does not allow to pass structs as Non-Type Template Parameters (NTTPs),
we cannot easily pass the policy around to other functions,
so it will be converted to legacy agent policies in many places.
Those are just structs with static data members holding the policy value as a type,
so they can be passed to templates.
Agent policy structs have historically been part of the public API, and should be removed in the future.
.. warning::
The kernel gets compiled for each target architecture (N many) that was provided to the compiler.
During each device pass, ``current_policy`` may return a different policy.
During the host pass, version A (runtime policy) compiles a single instantiation of the dispatch logic for all target architectures.
Version B (using ``dispatch_compute_cap``) compiles the dispatch logic for each distinct tuning policy (M many).
If we passed the selected tuning policy instead of the policy selector as a kernel template parameter,
the kernel template instantiation would be different for each tuning policy value and
we would compile O(M*N) kernels for version B instead of O(N).
Many kernels are short, since the functionality is extracted into the agent layer.
All the kernel does is derive the proper policy,
unwrap it to initialize the agent and call one of its ``Consume`` / ``Process`` functions.
Agents hold kernel bodies and are intended to be reused across multiple device-scope algorithms.
However, this is not a requirement and some kernels just contain their entire implementation themselves.
The default policy selector
------------------------------------
CUB contains a default policy selector for each dispatch function.
Because many dispatch functions are also used by CCCL.C,
which compiles them without proper type information,
we have to provide them in two forms,
a typeless and a typeful version.
.. code-block:: c++
struct AlgorithmPolicy { ... }; // unrelated to AgentPolicy
namespace detail::algorithm {
struct policy_selector {
type_t accum_t;
op_kind_t operation_t;
int offset_size;
int accum_size;
constexpr auto operator()(::cuda::compute_capability cc) const -> AlgorithmPolicy {
// parameter selection across target compute capabilities and input characteristics (possibly HUGE logic)
}
};
template <typename AccumT, typename OffsetT, typename ReductionOpT>
struct policy_selector_from_types {
constexpr auto operator()(cuda::compute_capability cc) const -> AlgorithmPolicy {
constexpr auto ps = policy_selector{
classify_type<AccumT>(),
classify_op<ReductionOpT>(),
int{sizeof(OffsetT)},
int{sizeof(AccumT)}
};
return ps(cc);
}
};
}
The ``policy_selector`` is intended to be used without template-parameter-based type information
and is thus suitable to be used in CCCL.C without a JIT compiler for host code.
It contains the necessary information on the algorithm's input as data members.
This policy is only used in the host code of the dispatch function.
Because it is not stateless anymore, CCCL.C overrides the kernel launcher used by CUB,
providing the kernel from a JIT-compiled instantiation that uses a proper stateless policy selector.
The kernel, and the dispatch function when not called from CCCL.C,
will use the second version of the policy selector, ``policy_selector_from_types``,
which offers proper template parameters to pass type information.
This type is stateless and delegates to a ``constexpr`` instance of the ``policy_selector``
with compile-time values derived from the template parameters.
The policy selection logic is thus the same for CCCL.C and CUB,
the type information is just provided differently.
Each dispatch function also has an associated concept for the policy selector it expects:
.. code-block:: c++
template <typename T, typename Policy>
concept policy_selector = requires(T pol_sel, cuda::compute_capability cc) {
requires std::regular<Policy>;
{ pol_sel(cc) } -> std::same_as<Policy>;
};
namespace detail::algorithm {
template <typename T>
concept algorithm_policy_selector = policy_selector<T, AlgorithmPolicy>;
}
The concept basically checks whether the policy selector can be called with a ``cuda::compute_capability``
and returns the expected policy struct.
The default policy selectors and related concept are defined in ``cub/device/dispatch/tuning/tuning_<algorithm>.cuh``.
Backward compatibility
------------------------------------
Legacy public dispatchers (e.g. ``DispatchReduce``) are deprecated.
They continue to work by translating the ``PolicyHub`` template parameter to the new ``policy_selector``
via a ``policy_selector_from_hub`` adapter:
.. code-block:: c++
template <typename PolicyHub>
struct policy_selector_from_hub {
constexpr auto operator()(cuda::compute_capability cc) const -> AlgorithmPolicy {
// conversion logic
}
};
This allows existing user code that passes custom policy hubs to dispatchers to continue working.
The legacy dispatchers, policy hubs and related logic are scheduled for removal in CCCL 4.0.
Policies
====================================
Policies describe the configuration of agents or just kernels with respect to performance.
They must not change functional behavior, but affect how work is mapped to the hardware
by defining certain parameters (items per thread, block size, etc.),
or choosing between algorithms.
Policies must be plain semiregular aggregates to allow using them during constant evaluation,
and with designated initializers:
.. code-block:: c++
struct AlgorithmPolicy {
int threads_per_block;
int items_per_thread;
cub::BlockLoadAlgorithm load_algorithm;
cub::CacheLoadModifier load_modifier;
friend constexpr bool operator==(const AlgorithmPolicy& lhs, const AlgorithmPolicy& rhs) { ... }
friend constexpr bool operator!=(const AlgorithmPolicy& lhs, const AlgorithmPolicy& rhs) { ... }
friend std::ostream& operator<<(std::ostream& os, const AlgorithmPolicy& p) { ... }
};
Tuning policies can have various complexities and contain nested structures.
Policies are defined in ``cub/device/dispatch/tuning/tuning_<algorithm>.cuh``.
Tunings
====================================
Because the values to parameterize an agent may vary a lot for different compile-time parameters,
the selection of values can involve complex logic.
Often, such tunings are found by experimentation or heuristic search.
See also :ref:`cub-tuning-infra`.
Tunings are expressed as logic and values inside the ``constexpr operator()`` of a policy selector.
Because of the complexity of some policy selectors, nested functions may be used.
Many policy selectors also implement a fallback logic,
where they try to find a matching tuning based on the input characteristics (policy selector data members),
but if no match is found, they fall back to an older target compute capability.
Here is an example:
.. code-block:: c++
struct sm90_tuning_values {
int items;
int threads;
int items_per_vec_load;
};
constexpr auto get_sm90_tuning(type_t accum_t, op_kind_t op, int offset_size, int accum_size)
-> std::optional<sm90_tuning_values> {
// provide tunings for certain operations, accumulator or offset types
if (op == op_kind_t::plus) {
if (accum_t == type_t::float32 && offset_size == 4 && accum_size == 4)
return sm90_tuning_values{16, 512, 2};
if (offset_size == 4 && accum_size == 8)
return sm90_tuning_values{15, 512, 2};
}
return {}; // no tuning available, causes fallback
}
struct policy_selector {
type_t accum_t;
op_kind_t operation_t;
int offset_size;
int accum_size;
constexpr auto operator()(cuda::compute_capability cc) const -> reduce_policy {
if (cc >= cuda::compute_capability{9, 0}) {
if (auto tuning = get_sm90_tuning(accum_t, operation_t, offset_size, accum_size)) {
return *tuning; // found a tuning, use it
}
// fall through to sm_80 if no matching tuning found
}
if (cc >= cuda::compute_capability{8, 0}) {
return { /* sm_80 default policy */ };
}
return { /* default policy for everything else */ };
}
};
In general, tunings are not exhaustive and usually only apply for specific combinations
of parameter values and a single compute capability.
This is because they originate from tuning benchmarks running for specific workloads on specific target architectures.
Generic fallbacks are often just retained from earlier days of CUB to not risk regressions,
or are based on heuristics trying to provide reasonable performance based on a model of the GPU architecture or algorithm.
Tunings for CUB algorithms reside in ``cub/device/dispatch/tuning/tuning_<algorithm>.cuh``.
Temporary storage usage
====================================
It's safe to reuse storage in the stream order:
.. code-block:: c++
cub::DeviceReduce::Sum(nullptr, storage_bytes, d_in, d_out, num_items, stream_1);
// allocate temp storage
cub::DeviceReduce::Sum(d_storage, storage_bytes, d_in, d_out, num_items, stream_1);
// fine not to synchronize stream
cub::DeviceReduce::Sum(d_storage, storage_bytes, d_in, d_out, num_items, stream_1);
// illegal, should call cudaStreamSynchronize(stream)
cub::DeviceReduce::Sum(d_storage, storage_bytes, d_in, d_out, num_items, stream_2);
Temporary storage management
====================================
Often times temporary storage for device-scope algorithms has a complex structure.
To simplify temporary storage management and make it safer,
we introduced ``cub::detail::temporary_storage::layout``:
.. code-block:: c++
cub::detail::temporary_storage::layout<2> storage_layout;
auto slot_1 = storage_layout.get_slot(0);
auto slot_2 = storage_layout.get_slot(1);
auto allocation_1 = slot_1->create_alias<int>();
auto allocation_2 = slot_1->create_alias<double>(42);
auto allocation_3 = slot_2->create_alias<char>(12);
if (condition)
{
allocation_1.grow(num_items);
}
if (d_temp_storage == nullptr)
{
temp_storage_bytes = storage_layout.get_size();
return;
}
storage_layout.map_to_buffer(d_temp_storage, temp_storage_bytes);
// different slots, safe to use simultaneously
use(allocation_1.get(), allocation_3.get(), stream);
// `allocation_2` alias `allocation_1`, safe to use in stream order
use(allocation_2.get(), stream);

17
cccl_upstream/docs/cub/developer/nvtx.rst Normal file
View File

@@ -0,0 +1,17 @@
.. _cub-developer-guide-nvtx:
NVTX
=====
The `NVIDIA Tools Extension SDK (NVTX) <https://nvidia.github.io/NVTX/>`_ is a cross-platform API
for annotating source code to provide contextual information to developer tools.
All device-scope algorithms in CUB are annotated with NVTX ranges,
allowing their start and stop to be visualized in profilers
like `NVIDIA Nsight Systems <https://developer.nvidia.com/nsight-systems>`_.
Only the public APIs available in the ``<cub/device/device_xxx.cuh>`` headers are annotated,
excluding direct calls to the dispatch layer.
NVTX annotations can be disabled by defining ``NVTX_DISABLE`` during compilation.
When CUB device algorithms are called on a stream subject to
`graph capture <https://developer.nvidia.com/blog/cuda-graphs/>`_,
the NVTX range is reported for the duration of capture (where no execution happens),
and not when a captured graph is executed later (the actual execution).

328
cccl_upstream/docs/cub/developer/test_overview.rst Normal file
View File

@@ -0,0 +1,328 @@
CUB Tests
###########################
.. warning::
CUB is in the progress of migrating to [Catch2](https://github.com/catchorg/Catch2) framework.
CUB tests rely on `CPM <https://github.com/cpm-cmake/CPM.cmake>`_ to fetch
`Catch2 <https://github.com/catchorg/Catch2>`_ that's used as our main testing framework.
Currently,
legacy tests coexist with Catch2 ones.
This guide is focused on new tests.
.. important::
Instead of including ``<catch2/catch.hpp>`` directly, use ``catch2_test_helper.h``.
.. code-block:: c++
#include <cub/block/block_scan.cuh>
#include <c2h/vector.h>
#include <c2h/catch2_test_helper.h>
Directory and File Naming
*************************************
Our tests can be found in the ``test`` directory.
Legacy tests have the following naming scheme: ``test_SCOPE_FACILITY.cu``.
For instance, here are the reduce tests:
.. code-block:: c++
test/test_warp_reduce.cu
test/test_block_reduce.cu
test/test_device_reduce.cu
Catch2-based tests have a different naming scheme: ``catch2_test_SCOPE_FACILITY.cu``.
The prefix is essential since that's how CMake finds tests
and distinguishes new tests from legacy ones.
Test Structure
*************************************
Base case
=====================================
Let's start with a simple example.
Say there's no need to cover many types with your test.
.. code-block:: c++
// 0) Define test name and tags
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]")
{
using type = std::int32_t;
constexpr int threads_per_block = 256;
constexpr int num_items = threads_per_block;
// 1) Allocate device input
c2h::device_vector<type> d_input(num_items);
// 2) Generate 3 random input arrays using Catch2 helper
c2h::gen(C2H_SEED(3), d_input);
// 3) Allocate output array
c2h::device_vector<type> d_output(d_input.size());
// 4) Copy device input to host
c2h::host_vector<key_t> h_reference = d_input;
// 5) Compute reference output
std::ALGORITHM(
thrust::raw_pointer_cast(h_reference.data()),
thrust::raw_pointer_cast(h_reference.data()) + h_reference.size());
// 6) Compute CUB output
SCOPE_ALGORITHM<threads_per_block>(d_input.data(),
d_output.data(),
d_input.size());
// 7) Compare device and host results
REQUIRE( d_input == d_output );
}
We introduce test cases with the ``C2H_TEST`` macro in (0).
This macro always takes two string arguments - a free-form test name and
one or more tags. Then, in (1), we allocate device memory using ``c2h::device_vector``.
``c2h::device_vector`` and ``c2h::host_vector`` behave similarly to their Thrust counterparts,
but are modified to provide more stable behavior in some testing edge cases.
.. important::
Always use ``c2h::host_vector<T>``/``c2h::device_vector<T>``
instead of ``thrust::host_vector<T>``/``thrust::device_vector<T>``,
unless the test code is being used for documentation examples.
Similarly, any thrust algorithms that executed on the device must be invoked with the
`c2h::device_policy` execution policy (not shown here) to support the same edge cases.
The memory is filled with random data in (2).
Generator ``c2h::gen`` takes at least two parameters.
The first one is a random generator seed.
Instead of providing a single value, we use the ``C2H_SEED`` macro.
The macro expects a number of seeds that has to be generated.
In the example above, we require three random seeds to be generated.
This leads to the whole test being executed three times
with different seed values.
Later, in (3), we allocate device output and host reference.
In (4), we allocate and populate the host input data.
Then, we perform the reference computation on the host in (5).
.. important::
Standard library algorithms (``std::``) have to be used where possible when computing reference solutions.
Afterwards, we launch the corresponding CUB algorithm in (6).
At this point, we have a reference solution on the CPU and a CUB solution on the GPU.
The two can be compared using Catch2's ``REQUIRE`` macro, which stops execution upon failure (preferred).
Catch2 also offers the ``CHECK`` macro, which continues test execution if the check fails.
If your test has to cover floating point types,
it's sufficient to replace ``REQUIRE( a == b )`` with ``REQUIRE_APPROX_EQ(a, b)``.
.. important::
Using ``c2h::gen`` for producing input data is strongly advised.
Do not use ``assert`` in tests, which is usually only enabled in Debug mode,
and we run CUB tests in Release mode.
If a custom (non-fundamental) type has to be tested, the following helper class template should be used:
.. code-block:: c++
using type = c2h::custom_type_t<c2h::accumulateable_t,
c2h::equal_comparable_t>;
Here we enumerate all the type properties that we are interested in.
The produced type ends up having ``operator==`` (from ``equal_comparable_t``)
and ``operator+`` (from ``accumulateable_t``).
More properties are available.
If a property is missing, please add it to the existing set in ``c2h``
instead of writing a custom type from scratch.
Generators
=====================================
We often need to test CUB algorithms against different inputs or problem sizes.
If these are **runtime values**, we can use the Catch2 ``GENERATE`` macro:
.. code-block:: c++
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]")
{
int num_items = GENERATE(1, 100, 1'000'000); // 0) Init. a variable with a generator
// ...
}
This will lead to the test being executed three times, once for each argument to ``GENERATE(...)``.
Multiple generators in a test inside the same scope will form the cartesian product of all combinations.
Please consult the `Catch2 documentation <https://github.com/catchorg/Catch2/blob/devel/docs/generators.md>`_
for more details.
``C2H_SEED(3)`` uses a generator expression internally.
Type Lists
=====================================
Since CUB is a generic library,
it's often required to test CUB algorithms against many types.
To do so,
it's sufficient to define a type list and provide it to the ``C2H_TEST`` macro.
This is useful for **compile-time** parameterization of tests.
.. code-block:: c++
// 0) Define type list
using types = c2h::type_list<std::uint8_t, std::int32_t>;
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]",
types) // 1) Provide it to the test case
{
// 2) Access current type with `c2h::get`
using type = typename c2h::get<0, TestType>;
// ...
}
This will lead to the test being compiled (instantiated) and run twice.
The first run will cause ``type`` to be ``std::uint8_t``.
The second one will cause ``type`` to be ``std::uint32_t``.
.. warning::
It's important to use types from the ``<cstdint>`` header
instead of built-in types like ``char`` and ``int``.
Multidimensional Configuration Spaces
=====================================
In most cases, the input data type is not the only compile-time parameter we want to vary.
For instance, you might need to test a block algorithm for different data types
**and** different thread block sizes.
To do so, you can add another type list as follows:
.. code-block:: c++
using block_sizes = c2h::enum_type_list<int, 128, 256>;
using types = c2h::type_list<std::uint8_t, std::int32_t>;
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]",
types, block_sizes)
{
using type = typename c2h::get<0, TestType>;
constexpr int threads_per_block = c2h::get<1, TestType>::value;
// ...
}
The code above leads to the following combinations being compiled:
- ``type = std::uint8_t``, ``threads_per_block = 128``
- ``type = std::uint8_t``, ``threads_per_block = 256``
- ``type = std::int32_t``, ``threads_per_block = 128``
- ``type = std::int32_t``, ``threads_per_block = 256``
As an example, the following test case includes both multidimensional configuration spaces
and multiple random sequence generations.
.. code-block:: c++
using block_sizes = c2h::enum_type_list<int, 128, 256>;
using types = c2h::type_list<std::uint8_t, std::int32_t>;
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]",
types, block_sizes)
{
using type = typename c2h::get<0, TestType>;
constexpr int threads_per_block = c2h::get<1, TestType>::value;
// ...
c2h::device_vector<type> d_input(5);
c2h::gen(C2H_SEED(2), d_input);
}
The code above leads to the following combinations being compiled:
- ``type = std::uint8_t``, ``threads_per_block = 128``, 1st random generated input sequence
- ``type = std::uint8_t``, ``threads_per_block = 256``, 1st random generated input sequence
- ``type = std::int32_t``, ``threads_per_block = 128``, 1st random generated input sequence
- ``type = std::int32_t``, ``threads_per_block = 256``, 1st random generated input sequence
- ``type = std::uint8_t``, ``threads_per_block = 128``, 2nd random generated input sequence
- ``type = std::uint8_t``, ``threads_per_block = 256``, 2nd random generated input sequence
- ``type = std::int32_t``, ``threads_per_block = 128``, 2nd random generated input sequence
- ``type = std::int32_t``, ``threads_per_block = 256``, 2nd random generated input sequence
Each new generator multiplies the number of execution times by its number of seeds. That means
that if there were further more sequence generators (``c2h::gen(C2H_SEED(X), ...)``) on the
example above the test would execute X more times and so on.
Speedup Compilation Time
=====================================
Since type lists in the ``C2H_TEST`` form a Cartesian product,
compilation time grows quickly with every new dimension.
To keep the compilation process parallelized,
it's possible to rely on our ``%PARAM%`` machinery:
.. code-block:: c++
// %PARAM% BLOCK_SIZE bs 128:256
using block_sizes = c2h::enum_type_list<int, BLOCK_SIZE>;
using types = c2h::type_list<std::uint8_t, std::int32_t>;
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]",
types, block_sizes)
{
using type = typename c2h::get<0, TestType>;
constexpr int threads_per_block = c2h::get<1, TestType>::value;
// ...
}
The comment with ``%PARAM%`` is recognized by our CMake scripts.
It leads to multiple executables being produced from a single test source.
.. code-block:: bash
bin/cub.test.scope_algorithm.bs_128
bin/cub.test.scope_algorithm.bs_256
Multiple ``%PARAM%`` comments can be specified forming another Cartesian product.
Final Test
=====================================
Let's consider the final test that illustrates all of the tools we discussed above:
.. code-block:: c++
// %PARAM% BLOCK_SIZE bs 128:256
using block_sizes = c2h::enum_type_list<int, BLOCK_SIZE>;
using types = c2h::type_list<std::uint8_t, std::int32_t>;
C2H_TEST("SCOPE FACILITY works with CONDITION", "[FACILITY][SCOPE]",
types, block_sizes)
{
using type = typename c2h::get<0, TestType>;
constexpr int threads_per_block = c2h::get<1, TestType>::value;
constexpr int max_num_items = threads_per_block;
c2h::device_vector<type> d_input(
GENERATE_COPY(take(2, random(0, max_num_items))));
c2h::gen(C2H_SEED(3), d_input);
c2h::device_vector<type> d_output(d_input.size());
SCOPE_ALGORITHM<threads_per_block>(d_input.data(),
d_output.data(),
d_input.size());
REQUIRE( d_input == d_output );
const type expected_sum = 4;
const type sum = thrust::reduce(c2h::device_policy, d_output.cbegin(), d_output.cend());
REQUIRE( sum == expected_sum);
}
Apart from discussed tools, here we also rely on ``Catch2`` to generate random input sizes
in the range ``[0, max_num_items]`` for our input vector ``d_input``.
Overall, the test will produce two executables.
Each of these executables is going to generate ``2`` input problem sizes.
For each problem size, ``3`` random vectors are generated.
As a result, we have ``12`` different tests.
The code also demonstrates the syntax and usage of ``c2h::device_policy`` with a Thrust algorithm.

24
cccl_upstream/docs/cub/developer/thread_level.rst Normal file
View File

@@ -0,0 +1,24 @@
.. _cub-developer-guide-thread-level:
Thread-level
*************
In contrast to algorithms at the warp/block/device layer,
single threaded functionality like ``cub::ThreadReduce``
is typically implemented as a sequential function and rarely exposed to the user.
.. code-block:: c++
template <
int LENGTH,
typename T,
typename ReductionOp,
typename PrefixT,
typename AccumT = detail::accumulator_t<ReductionOp, PrefixT, T>>
__device__ __forceinline__ AccumT ThreadReduce(
T (&input)[LENGTH],
ReductionOp reduction_op,
PrefixT prefix)
{
return ...;
}

157
cccl_upstream/docs/cub/developer/warp_level.rst Normal file
View File

@@ -0,0 +1,157 @@
.. _cub-developer-guide-warp-level:
Warp-level
************************************
CUB warp-level algorithms are specialized for execution by threads in the same CUDA warp.
These algorithms may only be invoked by ``1 <= n <= 32`` *consecutive* threads in the same warp.
Overview
====================================
Warp-level functionality is provided by types (classes) to provide encapsulation and enable partial template specialization.
For example, :cpp:struct:`cub::WarpReduce` is a class template:
.. code-block:: c++
template <typename T,
int LOGICAL_WARP_THREADS = 32>
class WarpReduce {
// ...
// (1) define `_TempStorage` type
// ...
_TempStorage &temp_storage;
public:
// (2) wrap `_TempStorage` in uninitialized memory
struct TempStorage : Uninitialized<_TempStorage> {};
__device__ __forceinline__ WarpReduce(TempStorage &temp_storage)
// (3) reinterpret cast
: temp_storage(temp_storage.Alias())
{}
// (4) actual algorithms
__device__ __forceinline__ T Sum(T input);
};
In CUDA, the hardware warp size is 32 threads.
However, CUB enables warp-level algorithms on "logical" warps of ``1 <= n <= 32`` threads.
The size of the logical warp is required at compile time via the ``LOGICAL_WARP_THREADS`` non-type template parameter.
This value is defaulted to the hardware warp size of ``32``.
There is a vital difference in the behavior of warp-level algorithms that depends on the value of ``LOGICAL_WARP_THREADS``:
- If ``LOGICAL_WARP_THREADS`` is a power of two - warp is partitioned into *sub*-warps,
each reducing its data independently from other *sub*-warps.
The terminology used in CUB: ``32`` threads are called hardware warp.
Groups with less than ``32`` threads are called *logical* or *virtual* warp since it doesn't correspond directly to any hardware unit.
- If ``LOGICAL_WARP_THREADS`` is **not** a power of two - there's no partitioning.
That is, only the first logical warp executes algorithm.
.. TODO: Add diagram showing non-power of two logical warps.
Temporary storage usage
====================================
Warp-level algorithms require temporary storage for scratch space and inter-thread communication.
The temporary storage needed for a given instantiation of an algorithm is known at compile time
and is exposed through the ``TempStorage`` member type definition.
It is the caller's responsibility to create this temporary storage and provide it to the constructor of the algorithm type.
It is possible to reuse the same temporary storage for different algorithm invocations,
but it is unsafe to do so without first synchronizing to ensure the first invocation is complete.
.. TODO: Add more explanation of the `TempStorage` type and the `Uninitialized` wrapper.
.. TODO: Explain if `TempStorage` is required to be shared memory or not.
.. code-block:: c++
using WarpReduce = cub::WarpReduce<int>;
// Allocate WarpReduce shared memory for four warps
__shared__ WarpReduce::TempStorage temp_storage[4];
// Get this thread's warp id
int warp_id = threadIdx.x / 32;
int aggregate_1 = WarpReduce(temp_storage[warp_id]).Sum(thread_data_1);
// illegal, has to add `__syncwarp()` between the two
int aggregate_2 = WarpReduce(temp_storage[warp_id]).Sum(thread_data_2);
// illegal, has to add `__syncwarp()` between the two
foo(temp_storage[warp_id]);
Specialization
====================================
The goal of CUB is to provide users with algorithms that abstract the complexities of achieving speed-of-light performance across a variety of use cases and hardware.
It is a CUB developer's job to abstract this complexity from the user by providing a uniform interface that statically dispatches to the optimal code path.
This is usually accomplished via customizing the implementation based on compile time information like the logical warp size, the data type, and the target architecture.
For example, :cpp:struct:`cub::WarpReduce` dispatches to two different implementations based on if the logical warp size is a power of two (described above):
.. code-block:: c++
using InternalWarpReduce = cuda::std::conditional_t<
IS_POW_OF_TWO,
detail::WarpReduceShfl<T, LOGICAL_WARP_THREADS>, // shuffle-based implementation
detail::WarpReduceSmem<T, LOGICAL_WARP_THREADS>>; // smem-based implementation
Specializations provide different shared memory requirements,
so the actual ``_TempStorage`` type is defined as:
.. code-block:: c++
using _TempStorage = typename InternalWarpReduce::TempStorage;
and algorithm implementation look like:
.. code-block:: c++
__device__ __forceinline__ T Sum(T input, int valid_items) {
return InternalWarpReduce(temp_storage)
.Reduce(input, valid_items, ::cuda::std::plus<>{});
}
``__CUDA_ARCH__`` cannot be used because it is conflicting with the PTX dispatch refactoring and limited NVHPC support.
Due to this limitation, we can't specialize on the PTX version.
``NV_IF_TARGET`` shall be used by specializations instead:
.. code-block:: c++
template <typename T, int LOGICAL_WARP_THREADS>
struct WarpReduceShfl
{
template <typename ReductionOp>
__device__ __forceinline__ T ReduceImpl(T input, int valid_items,
ReductionOp reduction_op)
{
// ... base case (SM < 80) ...
}
template <class U = T>
__device__ __forceinline__
typename std::enable_if<std::is_same_v<int, U> ||
std::is_same_v<unsigned int, U>,
T>::type
ReduceImpl(T input,
int, // valid_items
::cuda::std::plus<>) // reduction_op
{
T output = input;
NV_IF_TARGET(NV_PROVIDES_SM_80,
(output = __reduce_add_sync(member_mask, input);),
(output = ReduceImpl<::cuda::std::plus<>>(
input, LOGICAL_WARP_THREADS, ::cuda::std::plus<>{});));
return output;
}
};
Specializations are stored in the ``cub/warp/specializations`` directory.

138
cccl_upstream/docs/cub/developer_overview.rst Normal file
View File

@@ -0,0 +1,138 @@
Developer Overview
##########################
.. toctree::
:hidden:
:maxdepth: 2
developer/thread_level
developer/warp_level
developer/block_scope
developer/device_scope
developer/nvtx
developer/test_overview
This living document serves as a guide to the design of the internal structure of CUB.
CUB provides layered algorithms that correspond to the thread/warp/block/device hierarchy of threads in CUDA.
There are distinct algorithms for each layer and higher-level layers build on top of those below.
For example, CUB has four flavors of ``reduce``,
one for each layer: ``ThreadReduce, WarpReduce, BlockReduce``, and ``DeviceReduce``.
Each is unique in how it is invoked,
how many threads participate,
and on which thread(s) the result is valid.
These layers naturally build on each other.
For example, :cpp:struct:`cub::WarpReduce` uses :cpp:func:`cub::ThreadReduce`,
:cpp:struct:`cub::BlockReduce` uses :cpp:struct:`cub::WarpReduce`, etc.
:cpp:func:`cub::ThreadReduce`
- A normal function invoked and executed sequentially by a single thread that returns a valid result on that thread
- Single thread functions are usually an implementation detail and not exposed in CUB's public API
:cpp:struct:`cub::WarpReduce` and :cpp:struct:`cub::BlockReduce`
- A "cooperative" function where threads concurrently invoke the same function to execute parallel work
- The function's return value is well-defined only on the "first" thread (lowest thread index)
:cpp:struct:`cub::DeviceReduce`
- A normal function invoked by a single thread that spawns additional threads to execute parallel work
- Result is stored in the pointer provided to the function
- Function returns a ``cudaError_t`` error code
- Function does not synchronize the host with the device
The table below provides a summary of these functions:
.. list-table::
:class: table-no-stripes
:header-rows: 1
* - layer
- coop invocation
- parallel execution
- max threads
- valid result in
* - :cpp:func:`cub::ThreadReduce`
- :math:`-`
- :math:`-`
- :math:`1`
- invoking thread
* - :cpp:struct:`cub::WarpReduce`
- :math:`+`
- :math:`+`
- :math:`32`
- main thread
* - :cpp:struct:`cub::BlockReduce`
- :math:`+`
- :math:`+`
- :math:`1024`
- main thread
* - :cpp:struct:`cub::DeviceReduce`
- :math:`-`
- :math:`+`
- :math:`\infty`
- global memory
The details of how each of these layers are implemented is described below.
Common Patterns
************************************
While CUB's algorithms are unique at each layer,
there are commonalities among all of them:
- Algorithm interfaces are provided as *types* (classes)\ [1]_
- Algorithms need temporary storage
- Algorithms dispatch to specialized implementations depending on compile-time and runtime information
- Cooperative algorithms require the number of threads at compile time (template parameter)
Invoking any CUB algorithm follows the same general pattern:
#. Select the class for the desired algorithm
#. Query the temporary storage requirements
#. Allocate the temporary storage
#. Pass the temporary storage to the algorithm
#. Invoke it via the appropriate member function
An example of :cpp:struct:`cub::BlockReduce` demonstrates these patterns in practice:
.. code-block:: c++
__global__ void kernel(int* per_block_results)
{
// (1) Select the desired class
// `cub::BlockReduce` is a class template that must be instantiated for the
// input data type and the number of threads. Internally the class is
// specialized depending on the data type, number of threads, and hardware
// architecture. Type aliases are often used for convenience:
using BlockReduce = cub::BlockReduce<int, 128>;
// (2) Query the temporary storage
// The type and amount of temporary storage depends on the selected instantiation
using TempStorage = typename BlockReduce::TempStorage;
// (3) Allocate the temporary storage
__shared__ TempStorage temp_storage;
// (4) Pass the temporary storage
// Temporary storage is passed to the constructor of the `BlockReduce` class
BlockReduce block_reduce{temp_storage};
// (5) Invoke the algorithm
// The `Sum()` member function performs the sum reduction of `thread_data` across all 128 threads
int thread_data[4] = {1, 2, 3, 4};
int block_result = block_reduce.Sum(thread_data);
per_block_results[blockIdx.x] = block_result;
}
.. [1] Algorithm interfaces are provided as classes because it provides encapsulation for things like temporary storage requirements and enables partial template specialization for customizing an algorithm for specific data types or number of threads.
For more detailed descriptions of the respective algorithms levels see the individual sections below
- :ref:`thread-level algorithms<cub-developer-guide-thread-level>`
- :ref:`warp-level algorithms<cub-developer-guide-warp-level>`
- :ref:`block-scope algorithms<cub-developer-guide-block-scope>`
- :ref:`device-scope algorithms<cub-developer-guide-device-scope>`
There is additional information for :ref:`nvtx ranges <cub-developer-guide-nvtx>`

323
cccl_upstream/docs/cub/device_topk_requirements.rst Normal file
View File

@@ -0,0 +1,323 @@
:orphan:
.. _cub-topk-requirements:
Top-K: Determinism, Tie-Breaking, and Output Ordering
======================================================
This page describes how to control the result of the CUB top-k family of algorithms
(:cpp:struct:`cub::DeviceTopK` and :cpp:struct:`cub::DeviceBatchedTopK`) through the execution
environment. For :cpp:struct:`cub::DeviceBatchedTopK`, these requirements apply independently within
each segment. The same requirement model applies to every ``MaxKeys`` / ``MinKeys`` / ``MaxPairs`` /
``MinPairs`` entry point.
Two orthogonal concerns
-----------------------
Top-k algorithms answer two separate questions:
#. **Which items are returned?** (the result *set* / membership), controlled by
``cuda::execution::determinism`` and, when deterministic, optionally refined by
``cuda::execution::tie_break``.
#. **In what order are those items written to the output?** (the result *sequence*), controlled
independently by ``cuda::execution::output_ordering``.
Think of it this way: determinism (with an optional tie-break) first selects a *set* of *K* items.
Output ordering then arranges that fixed set into the output buffer. Changing output ordering never
changes *which* items are selected. Changing tie-breaking never dictates *how* equal-key items are
sequenced in the output (unless you also request a stable ordering, as described below).
**Determinism applies to set membership.** Even with a deterministic selection, the *positions* of
the selected items in the output buffer may still vary unless you also request a specific output
ordering. Non-determinism arises only when more elements compare equal at the selection boundary
than there are remaining slots in the top-*K*. For example, with *K* = 3 and four elements tied for
the third-largest position, the algorithm must choose three of the four, and that choice is the
source of variability.
**Output ordering applies to the result sequence.** Once the result set is fixed, output ordering
specifies how those *K* items are laid out in the output buffer.
.. _cub-topk-default-behavior:
Default behavior
----------------
When you do **not** specify any of these requirements, the top-k algorithms provide their strongest
reproducibility guarantees. The committed default contract is:
* ``cuda::execution::determinism::gpu_to_gpu`` for a deterministic result set,
* ``cuda::execution::tie_break::prefer_smaller_index`` to resolve ties at the selection boundary
toward the smaller (lower) source index,
* ``cuda::execution::output_ordering::stable_sorted`` to write output sorted by key, with equal
keys ordered by source index.
In other words, by default you get the same items, in the same positions, run after run and across
GPUs of the same architecture. You opt **out** of these guarantees (by requiring weaker properties
such as ``cuda::execution::determinism::not_guaranteed`` and
``cuda::execution::output_ordering::unsorted``) to obtain faster implementations.
``determinism`` and ``tie_break`` are coupled. You specify **both** of them (inside a single
``cuda::execution::require(...)``) or **neither** (to take the default). A specified ``tie_break`` of
``prefer_smaller_index`` or ``prefer_larger_index`` pins the result set across GPUs and therefore
requires ``determinism::gpu_to_gpu``. See :ref:`cub-topk-set-membership` for the full table.
.. note::
**Current support.** This initial API surface only implements the fully opted-out configuration.
For :cpp:struct:`cub::DeviceBatchedTopK` it must be requested **explicitly** as
``cuda::execution::require(cuda::execution::determinism::not_guaranteed,
cuda::execution::tie_break::unspecified, cuda::execution::output_ordering::unsorted)``
(:cpp:struct:`cub::DeviceTopK` has no tie-break dimension yet and omits the ``tie_break`` token).
The algorithms ``static_assert`` for any other combination (including an empty, no-requirement
environment), so the deterministic default described above cannot yet be exercised in code. The
deterministic, tie-broken, and (stable-)sorted modes documented here define the committed long-term
contract and will become available (including as the no-requirement default) as those code paths
land.
Requirements reference
----------------------
Determinism (``cuda::execution::determinism``)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
.. list-table::
:header-rows: 1
:widths: 25 75
* - Value
- Meaning
* - ``not_guaranteed``
- No reproducibility guarantee. Among tied elements at the selection boundary, any valid subset
may be returned. Enables the fastest implementations.
* - ``run_to_run``
- The result set is identical across repeated invocations on the same GPU with the same input.
The tie-breaking policy is implementation-defined. Pinning a specific tie-break is not
available at this level and requires ``gpu_to_gpu``.
* - ``gpu_to_gpu``
- The result set is identical across different GPUs of the same architecture. This is the only
level that may be combined with an explicit ``tie_break`` (``prefer_smaller_index`` or
``prefer_larger_index``), which then fully pins the result set for a given input.
Tie-break (``cuda::execution::tie_break``)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
A specified ``tie_break`` of ``prefer_smaller_index`` or ``prefer_larger_index`` pins the result set
across GPUs, so it requires ``determinism::gpu_to_gpu``. Pairing it with ``run_to_run`` or
``not_guaranteed`` is rejected at compile time. ``determinism`` and ``tie_break`` must always be
specified together (or both omitted to take the default). Use ``tie_break::unspecified`` to leave the
boundary policy to the implementation, for example alongside ``not_guaranteed`` or ``run_to_run``.
.. list-table::
:header-rows: 1
:widths: 30 70
* - Value
- Meaning
* - ``unspecified``
- Any deterministic tie-break is acceptable, and the implementation chooses. Valid with any
determinism level (including ``not_guaranteed`` and ``run_to_run``).
* - ``prefer_smaller_index`` *(default)*
- Among elements that compare equal at the boundary, prefer those with the **smaller** source
index. Requires ``determinism::gpu_to_gpu``.
* - ``prefer_larger_index``
- Among elements that compare equal at the boundary, prefer those with the **larger** source
index. Requires ``determinism::gpu_to_gpu``.
Output ordering (``cuda::execution::output_ordering``)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
.. list-table::
:header-rows: 1
:widths: 25 75
* - Value
- Meaning
* - ``unsorted``
- No guarantee on output order. The same result set may appear in different permutations across
runs.
* - ``sorted``
- Output is sorted by key value (descending for ``Max*``, ascending for ``Min*``). Among
elements with equal keys, the relative order is **unspecified**.
* - ``stable_sorted``
- Output is sorted by key value, and among equal keys the relative order matches the **input
order** (smaller source index first). With a fully pinned result set (an explicit
``tie_break``) this fully determines the output, so the result is bit-identical even across
GPUs of the same architecture.
Composing requirements
----------------------
Requirements compose into a single ``cuda::execution::require(...)`` argument, which is placed in the
execution environment alongside other properties such as a stream:
.. code-block:: c++
auto env = cuda::std::execution::env{
cuda::execution::require(
cuda::execution::determinism::gpu_to_gpu,
cuda::execution::tie_break::prefer_smaller_index,
cuda::execution::output_ordering::sorted),
stream_ref};
.. _cub-topk-set-membership:
Which items are selected?
-------------------------
Determinism and tie-break together control **set membership**. They are always specified as a pair
(or both omitted to take the default). Rows below are the ``determinism`` requirement and columns are
the paired ``tie_break`` requirement. Cells marked *(compile error)* are rejected by a ``static_assert``.
.. list-table::
:header-rows: 1
:stub-columns: 1
:widths: 22 26 26 26
* - ``determinism``
- ``tie_break::unspecified``
- ``tie_break::prefer_smaller_index``
- ``tie_break::prefer_larger_index``
* - ``not_guaranteed``
- Non-deterministic (fast path)
- *(compile error)*
- *(compile error)*
* - ``run_to_run``
- Deterministic, implementation-defined tie-break
- *(compile error)*
- *(compile error)*
* - ``gpu_to_gpu``
- Deterministic, implementation-defined tie-break
- Deterministic, ties toward the **smaller** source index
- Deterministic, ties toward the **larger** source index
Reading the table:
* A specified ``tie_break`` of ``prefer_smaller_index`` or ``prefer_larger_index`` pins the result set
across GPUs, which is a ``gpu_to_gpu`` guarantee. Requesting it alongside ``not_guaranteed`` or
``run_to_run`` is a compile error, because you must acknowledge the ``gpu_to_gpu`` determinism you
receive.
* With ``tie_break::unspecified`` the implementation chooses the boundary policy. ``run_to_run`` and
``gpu_to_gpu`` then differ only in *scope*: identical results on the same GPU versus across GPUs of
the same architecture.
* Omitting **both** requirements selects the default (``gpu_to_gpu`` with ``prefer_smaller_index``),
which is the bottom-middle cell.
.. note::
This determinism and tie_break pairing rule is currently enforced only by
:cpp:struct:`cub::DeviceBatchedTopK`. :cpp:struct:`cub::DeviceTopK` does not yet inspect
``tie_break``, so it still accepts requirement combinations that ``cub::DeviceBatchedTopK`` rejects.
The same enforcement will be added to ``cub::DeviceTopK`` in the next major release of CCCL (4.0).
Worked example: set membership x output ordering
-------------------------------------------------
Consider ``cub::DeviceTopK::MaxKeys`` with *K* = 3 on this input:
.. code-block:: text
index : 0 1 2 3 4 5
key : 10 8 8 8 6 5
The top three keys are ``10`` and two ``8``\ s. Four elements compare equal at the boundary (the
``8``\ s at indices 1, 2, 3), but only two can be kept. That is the tie. The notation ``key@index``
identifies an element by both its key and its source position (for example ``8@2`` is the ``8`` at
index 2).
The table below shows **two runs on the same input** for each combination. Compare the two runs
within a cell to see whether the output order varies. Compare across rows to see whether the set
membership varies.
.. list-table::
:header-rows: 1
:widths: 28 24 24 24
* - ``require(...)``
- ``output_ordering::unsorted``
- ``output_ordering::sorted``
- ``output_ordering::stable_sorted``
* - ``determinism::not_guaranteed,``
``tie_break::unspecified``
- | Run 1: ``[8@2, 10@0, 8@1]``
| Run 2: ``[8@3, 10@0, 8@1]``
| Different sets *and* orders
- | Run 1: ``[10@0, 8@2, 8@1]``
| Run 2: ``[10@0, 8@1, 8@3]``
| Different sets, sorted by key
- | Run 1: ``[10@0, 8@1, 8@2]``
| Run 2: ``[10@0, 8@1, 8@3]``
| Different sets, equal keys in input order
* - ``determinism::run_to_run,``
``tie_break::unspecified``
- | Run 1: ``[8@3, 10@0, 8@1]``
| Run 2: ``[10@0, 8@1, 8@3]``
| Same set ``{10@0, 8@1, 8@3}``, order may vary
- | Run 1: ``[10@0, 8@3, 8@1]``
| Run 2: ``[10@0, 8@1, 8@3]``
| Same set, equal-key order unspecified
- | Run 1: ``[10@0, 8@1, 8@3]``
| Run 2: ``[10@0, 8@1, 8@3]``
| Same set, equal keys always in input order
* - ``determinism::gpu_to_gpu,``
``tie_break::prefer_smaller_index``
- | Run 1: ``[8@2, 10@0, 8@1]``
| Run 2: ``[10@0, 8@1, 8@2]``
| Same set ``{10@0, 8@1, 8@2}``, order may vary
- | Run 1: ``[10@0, 8@2, 8@1]``
| Run 2: ``[10@0, 8@1, 8@2]``
| Same set, equal-key order unspecified
- | Run 1: ``[10@0, 8@1, 8@2]``
| Run 2: ``[10@0, 8@1, 8@2]``
| Same set, equal keys always in input order
* - ``determinism::gpu_to_gpu,``
``tie_break::prefer_larger_index``
- | Run 1: ``[8@3, 10@0, 8@2]``
| Run 2: ``[10@0, 8@2, 8@3]``
| Same set ``{10@0, 8@2, 8@3}``, order may vary
- | Run 1: ``[10@0, 8@3, 8@2]``
| Run 2: ``[10@0, 8@2, 8@3]``
| Same set, equal-key order unspecified
- | Run 1: ``[10@0, 8@2, 8@3]``
| Run 2: ``[10@0, 8@2, 8@3]``
| Same set, equal keys always in input order
Reading the matrix:
.. list-table::
:header-rows: 1
:widths: 45 55
* - Observation
- Where to look
* - Set membership varies across runs
- ``not_guaranteed`` row: Run 1 keeps ``8@2``, Run 2 keeps ``8@3``
* - Set membership fixed, order varies
- ``run_to_run`` + ``unsorted``: both runs return ``{10@0, 8@1, 8@3}`` in different permutations
* - Set membership fixed, sorted but unstable among equal keys
- ``run_to_run`` + ``sorted``: both runs start with ``10@0``, but ``8@1`` and ``8@3`` may swap
* - Fully pinned: same set and same order
- ``gpu_to_gpu`` + ``tie_break::prefer_smaller_index`` + ``stable_sorted``: both runs yield
``[10@0, 8@1, 8@2]``
* - Tie-break changes the set, not just the order
- Compare ``prefer_smaller_index`` vs ``prefer_larger_index``: ``8@2`` vs ``8@3``
Choosing requirements
---------------------
.. list-table::
:header-rows: 1
:widths: 55 45
* - Goal
- Suggested ``require(...)``
* - Maximum performance, exact result unimportant
- ``determinism::not_guaranteed, tie_break::unspecified, output_ordering::unsorted``
* - Reproducible result set, order does not matter
- ``determinism::run_to_run, tie_break::unspecified, output_ordering::unsorted``
* - Reproducible result set with an explicit boundary policy
- ``determinism::gpu_to_gpu, tie_break::prefer_{smaller,larger}_index, output_ordering::unsorted``
* - Reproducible, key-sorted output
- the above + ``output_ordering::sorted``
* - Reproducible, key-sorted output with input-order stability among ties
- the above + ``output_ordering::stable_sorted`` (a fully pinned set plus stable-sorted output
is bit-identical, including across GPUs)

120
cccl_upstream/docs/cub/device_wide.rst Normal file
View File

@@ -0,0 +1,120 @@
.. _device-module:
Device-Wide Primitives
======================
.. toctree::
:glob:
:hidden:
:maxdepth: 2
api/device
Almost all of CUB's device-wide APIs come in two flavors:
* the traditional two-phase style that requires calling the API twice and managing temporary storage explicitly,
* and the newer single-phase style where temporary storage is obtained from a memory resource in the execution environment.
Some APIs that do not require any temporary storage may also have a traditional single-phase form in addition to the newer environment-based one.
.. _device-temp-storage:
Two-Phase API (explicit temporary storage management)
+++++++++++++++++++++++++++++++++++++++++++++++++++++
Traditional two-phase APIs can be recognized by taking ``void* d_temp_storage, size_t& temp_storage_bytes`` as their first two parameters.
They follow a two-phase usage pattern that requires three steps:
1. **Query Phase**: The algorithm is called the first time with ``d_temp_storage = nullptr`` to determine the required temporary storage size.
The required size is written to ``temp_storage_bytes`` without dereferencing iterators or launching kernels.
2. **Temporary storage allocation**: The user is responsible for allocating device-accessible memory of at least ``temp_storage_bytes`` bytes.
No special alignment is required.
3. **Execution Phase**: The algorithm is called the second time with ``d_temp_storage`` pointing to the allocated device memory, performing the actual operation.
In principle, the query phase and execution phase must call the same CUB API.
This means in detail:
* **Template arguments**: The query call must use the same template arguments as the execution call, so they share the same template instantiation.
* **Argument values**: Regarding function parameters, only the values of the ``d_temp_storage``, ``temp_storage_bytes``,
and problem-size related arguments (like number of elements, number of segments, segment sizes, etc.) may be read during the query phase.
No other parameters (like input/output iterators, initial values, etc.) are accessed during the query phase, so their values may be indeterminate.
During the query phase, the API will return before launching any kernels or touching user storage.
* **Current device**: The computed temporary storage size is valid only when the execution phase runs on the same current CUDA device as the query.
Re-run the query if the current device changes between phases.
Example pattern:
.. literalinclude:: ../../cub/examples/device/example_device_reduce.cu
:language: c++
:dedent:
:start-after: example-begin temp-storage-query
:end-before: example-end temp-storage-query
Environment API (single phase)
++++++++++++++++++++++++++++++
The environment-based API is available for all CUB device-wide algorithms.
They remove the split of query/execute phase and manually obtaining the temporary storage.
Instead, the temporary storage is automatically requested from a memory resource queried from the execution environment argument.
The environment supports further properties like passing a stream or an execution requirement in addition to a memory resource.
Key properties of the environment argument:
- It is a defaulted parameter and appears as the last argument.
- Streams like `cudaStream_t` or `cuda::stream_ref` can be passed as environments directly, or added to the environment.
- You can select the memory resource (CCCL-provided or custom) used for internal allocations.
- Supported algorithms accept determinism requirements (for example, ``cuda::execution::determinism::gpu_to_gpu``).
- Multiple properties compose into a single centralized argument by wrapping them into a ``cuda::execution::env`` object.
Example pattern:
.. literalinclude:: ../../cub/examples/device/example_device_reduce_env.cu
:language: c++
:dedent:
:start-after: example-begin env-overload-setup
:end-before: example-end env-overload-setup
.. literalinclude:: ../../cub/examples/device/example_device_reduce_env.cu
:language: c++
:dedent:
:start-after: example-begin env-overload-run
:end-before: example-end env-overload-run
Further information on CUB execution environments can be found in
:ref:`Execution Environments <cub-environment>`.
API overview
++++++++++++
In the following, the various groups of CUB device-wide algorithms are listed,
linking to their respective documentation.
CUB device-level single-problem parallel algorithms:
* :cpp:struct:`cub::DeviceAdjacentDifference` computes the difference between adjacent elements residing within device-accessible memory
* :cpp:struct:`cub::DeviceFor` provides device-wide, parallel operations for iterating over data residing within device-accessible memory
* :cpp:struct:`cub::DeviceHistogram` constructs histograms from data samples residing within device-accessible memory
* :cpp:struct:`cub::DevicePartition` partitions data residing within device-accessible memory
* :cpp:struct:`cub::DeviceMerge` merges two sorted sequences in device-accessible memory into a single one
* :cpp:struct:`cub::DeviceMergeSort` sorts items residing within device-accessible memory
* :cpp:struct:`cub::DeviceRadixSort` sorts items residing within device-accessible memory using radix sorting method
* :cpp:struct:`cub::DeviceReduce` computes reduction of items residing within device-accessible memory
* :cpp:struct:`cub::DeviceRunLengthEncode` demarcating "runs" of same-valued items within a sequence residing within device-accessible memory
* :cpp:struct:`cub::DeviceScan` computes a prefix scan across a sequence of data items residing within device-accessible memory
* :cpp:struct:`cub::DeviceSelect` compacts data residing within device-accessible memory
* :cpp:struct:`cub::DeviceTransform` transforms elements from multiple input sequences into an output sequence
* :cpp:struct:`cub::DeviceTopK` finds the largest (or smallest) K items from an unordered list residing within device-accessible memory
CUB device-level segmented-problem (batched) parallel algorithms:
* :cpp:struct:`cub::DeviceSegmentedSort` computes batched sort across non-overlapping sequences of data residing within device-accessible memory
* :cpp:struct:`cub::DeviceSegmentedRadixSort` computes batched radix sort across non-overlapping sequences of data residing within device-accessible memory
* :cpp:struct:`cub::DeviceSegmentedReduce` computes reductions across multiple sequences of data residing within device-accessible memory
* :cpp:struct:`cub::DeviceSegmentedScan` computes prefix scans across multiple sequences of data residing within device-accessible memory
* :cpp:struct:`cub::DeviceBatchedTopK` finds the largest (or smallest) K items from each of multiple unordered lists (segments) residing within device-accessible memory
* :cpp:struct:`cub::DeviceCopy` provides device-wide, parallel operations for batched copying of data residing within device-accessible memory
* :cpp:struct:`cub::DeviceMemcpy` provides device-wide, parallel operations for batched copying of data residing within device-accessible memory
* :cpp:struct:`cub::DeviceFind` provides vectorized binary search algorithms

253
cccl_upstream/docs/cub/environment.rst Normal file
View File

@@ -0,0 +1,253 @@
.. _cub-environment:
Execution Environments
======================
Most CUB device-wide algorithms accept an optional *execution environment* as their last
argument. The environment is a single object that bundles everything CUB needs to know
about *how* to run the algorithm.
Part of what it carries is familiar: the CUDA stream, which CUB algorithms have always
accepted, now simply travels inside the environment. The rest are controls that had no
place in the classic API and are enabled by environments:
:ref:`determinism requirements <cub-env-determinism>`,
:ref:`custom tuning policy selectors <cub-env-tuning>`, and
:ref:`memory resources <cub-env-memory-resource>`. All properties are optional and freely
composable with each other.
This page explains what the CUB environment APIs are, and how to use them.
.. contents::
:local:
:depth: 2
Why environments?
-----------------
The classic two-phase CUB API requires three steps: query temporary-storage size, allocate,
then execute. That is fine for situations where precise control over the temporary storage
allocation is required, or when the temporary storage size needs to be queried independently
of allocating it. For example, when the temporary storage is combined with other storage, e.g.
for an output, into a single allocation. But in many cases this is not needed and the two-step
API is just repetitive boilerplate.
The environment-based single-phase API collapses all of that into one call. The algorithm
queries the environment for the properties it needs — for example, the stream to run on,
or a memory resource to allocate temporary storage from — and then executes. Properties
an algorithm does not use are simply ignored, which is what makes one environment safe to
pass to many different algorithms:
.. code-block:: c++
cub::DeviceReduce::Sum(d_input, d_output, num_items, env);
.. note::
The environment argument is entirely
optional: since it is defaulted, an algorithm can be invoked with no temporary-storage
arguments and no environment at all, and every property falls back to its default
(see :ref:`Default behavior <cub-environment-fallback>`):
.. code-block:: c++
cub::DeviceReduce::Sum(d_input, d_output, num_items);
.. _cub-env-building:
Building an environment
-----------------------
.. list explicitly what types are valid environments
Some types are already valid environments on their own and can be passed directly to an
algorithm. The most common case is a stream: ``cuda::stream_ref`` (or a raw
``cudaStream_t``) passed as the last argument is treated as an environment containing
just that stream.
.. literalinclude:: ../../cub/test/catch2_test_device_reduce_env_api.cu
:language: c++
:dedent:
:start-after: example-begin reduce-env-stream
:end-before: example-end reduce-env-stream
When you need more than one property, combine them with ``cuda::std::execution::env``.
Properties can be listed in any order. The example below
composes a stream, a memory pool for temporary storage, and a determinism requirement
(the required declarations are provided by ``<cuda/execution>``, ``<cuda/stream>``,
``<cuda/memory_pool>``, and ``<cuda/devices>``):
.. literalinclude:: ../../cub/examples/device/example_device_reduce_env.cu
:language: c++
:dedent:
:start-after: example-begin env-overload-setup
:end-before: example-end env-overload-setup
.. literalinclude:: ../../cub/examples/device/example_device_reduce_env.cu
:language: c++
:dedent:
:start-after: example-begin env-overload-run
:end-before: example-end env-overload-run
The same ``env`` object can be passed to multiple algorithm calls without rebuilding it each
time, see :ref:`cub-environment-reuse`.
How to use environments
-----------------------
The controls below are what environments enable beyond the classic API. A few algorithms
additionally accept algorithm-specific controls — e.g. tie-breaking and output-ordering
for :ref:`cub::DeviceTopK <cub-topk-requirements>`. The set of supported properties
keeps growing.
.. _cub-env-determinism:
Determinism Requirements
~~~~~~~~~~~~~~~~~~~~~~~~
Use ``cuda::execution::require`` to request a guarantee, like a reproducible/deterministic execution:
.. literalinclude:: ../../cub/test/catch2_test_device_reduce_env_api.cu
:language: c++
:dedent:
:start-after: example-begin reduce-env-determinism
:end-before: example-end reduce-env-determinism
Multiple determinism levels are available. The meaning of each is described in the
:ref:`CCCL determinism overview <cccl-determinism>`. Which algorithms support which levels,
and each algorithm's default, are listed in the :ref:`CUB determinism support matrix
<cub-determinism>`. Requesting a level an algorithm does not support is rejected at
compile time.
.. _cub-env-tuning:
Custom Tuning Policy Selectors
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
Pass a custom policy selector through the environment to override CUB's built-in tuning.
A policy selector answers the question "which tuning parameters should this algorithm use
on this GPU?": it is a function object that CUB calls with the GPU's compute capability,
and that returns the tuning parameters — threads per block, items per thread, and so on —
packaged as the algorithm's policy (here ``cub::MergePolicy``):
.. literalinclude:: ../../cub/test/catch2_test_device_merge_env_api.cu
:language: c++
:dedent:
:start-after: example-begin merge-keys-policy-selector
:end-before: example-end merge-keys-policy-selector
The selector is wrapped with ``cuda::execution::tune`` and passed as (part of) the
environment:
.. literalinclude:: ../../cub/test/catch2_test_device_merge_env_api.cu
:language: c++
:dedent:
:start-after: example-begin merge-keys-tuning
:end-before: example-end merge-keys-tuning
.. seealso:: :ref:`cub-policy-selectors` - full guide on defining and composing policy
selectors, including the requirements a policy selector must satisfy.
.. _cub-env-memory-resource:
Memory Resources
~~~~~~~~~~~~~~~~
The memory resource controls where an algorithm's temporary storage is allocated from.
Memory resource types are valid environments on their own, so they can be passed directly
or composed with other properties:
.. code-block:: c++
auto pool = cuda::device_default_memory_pool(cuda::devices[0]);
// Pass directly...
cub::DeviceReduce::Sum(d_input, d_output, num_items, pool);
// ...or compose with other properties
auto env = cuda::std::execution::env{stream, pool};
cub::DeviceReduce::Sum(d_input, d_output, num_items, env);
Temporary storage is allocated from the memory resource on the algorithm's stream before
execution and released on the same stream afterwards. When no memory resource is present
in the environment, CUB falls back to a stream-ordered ``cudaMallocAsync``/``cudaFree`` allocator
(see :ref:`Default behavior <cub-environment-fallback>`).
.. TODO: Add Guarantees sub-section after #9278 is merged.
.. _cub-environment-reuse:
Reusing an environment across multiple calls
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
An ``env`` object can be built once and passed to as many algorithm
calls as you like:
.. code-block:: c++
auto stream = cuda::stream{cuda::devices[0]};
auto pool = cuda::device_default_memory_pool(cuda::devices[0]);
auto env = cuda::std::execution::env{cuda::stream_ref{stream}, pool};
cub::DeviceScan::ExclusiveSum(d_a, d_out_a, n, env);
cub::DeviceReduce::Sum(d_b, d_out_b, n, env);
cub::DeviceSelect::If(d_c, d_out_c, d_num_selected, n, my_predicate, env);
stream.sync();
All three calls share the same stream and memory pool. Temporary storage is allocated and
released from the pool independently for each call.
.. note::
When a tuning policy selector is embedded in the environment, it applies to *every*
algorithm that can be tuned by this selector. For example, a policy selector returning
a ``cub::ReducePolicy`` will be used by all calls to ``cub::DeviceReduce::*`` that use
the same environment. If two algorithms using the same environment need different
tunings, build separate environments.
.. _cub-environment-fallback:
Default behavior
----------------
The environment argument is optional on every algorithm that supports it. CUB applies the
following defaults when a property is absent from the environment (or when no environment
is passed at all):
.. list-table::
:header-rows: 1
:widths: 25 75
* - Property
- Default when absent
* - Stream
- The default CUDA stream (``cudaStream_t{}``).
* - Memory resource
- A stream-ordered allocator based on ``cudaMallocAsync``. Temporary storage is
allocated on the stream the algorithm runs on.
* - Determinism
- Algorithm-specific. See :ref:`cub-determinism`.
* - Policy selector
- CUB's built-in selector, returning architecture-tuned defaults for the current
device.
Missing properties never cause an error. The algorithm simply falls back to its default for
that property.
..
TODO(gonidelis): link to the developer-facing environments page (queries, CPOs, prop,
building custom environment types) once it lands via
https://github.com/NVIDIA/cccl/pull/10013
See also
--------
- :ref:`cub-determinism` - per-algorithm determinism support matrix
- :ref:`cub-policy-selectors` - defining and composing custom policy selectors
- :ref:`device-module` - overview of the device-wide API and the two-phase alternative
- :ref:`cccl-determinism` - CCCL-level determinism concepts

487
cccl_upstream/docs/cub/index.rst Normal file
View File

@@ -0,0 +1,487 @@
.. _cub-module:
CUB
==================================================
.. toctree::
:hidden:
:maxdepth: 3
Overview <self>
thread_level
warp_wide
block_wide
device_wide
environment
determinism
benchmarking
tuning
tuning_infra
API reference <api/index>
developer_overview
What is CUB?
==================================================
CUB provides state-of-the-art, reusable software components for every layer
of the CUDA programming model:
* **Parallel primitives**
* :ref:`Thread <thread-module>` primitives
* Thread-level reduction, etc.
* Safely specialized for each underlying CUDA architecture
* :ref:`Warp-wide <warp-module>` "collective" primitives
* Cooperative warp-wide prefix scan, reduction, etc.
* Safely specialized for each underlying CUDA architecture
* :ref:`Block-wide <block-module>` "collective" primitives
* Cooperative I/O, sort, scan, reduction, histogram, etc.
* Compatible with arbitrary thread block sizes and types
* :ref:`Device-wide <device-module>` primitives
* Parallel sort, prefix scan, reduction, histogram, etc.
* Compatible with CUDA dynamic parallelism
* **Utilities**
* **Fancy iterators**
* **Thread and thread block I/O**
* **PTX intrinsics**
* **Device, kernel, and storage management**
.. _collective-primitives:
CUB's collective primitives
==================================================
Collective software primitives are essential for constructing high-performance,
maintainable CUDA kernel code. Collectives allow complex parallel code to be
re-used rather than re-implemented, and to be re-compiled rather than
hand-ported.
.. figure:: ../img/cub_overview.png
:align: center
:alt: Orientation of collective primitives within the CUDA software stack
:name: fig_cub_overview
Orientation of collective primitives within the CUDA software stack
As a SIMT programming model, CUDA engenders both **scalar** and
**collective** software interfaces. Traditional software
interfaces are *scalar* : a single thread invokes a library routine to perform some
operation (which may include spawning parallel subtasks). Alternatively, a *collective*
interface is entered simultaneously by a group of parallel threads to perform
some cooperative operation.
CUB's collective primitives are not bound to any particular width of parallelism
or data type. This flexibility makes them:
* **Adaptable** to fit the needs of the enclosing kernel computation
* **Trivially tunable** to different grain sizes (threads per block, items per thread, etc.)
Thus CUB is *CUDA Unbound*.
An example (block-wide sorting)
==================================================
The following code snippet presents a CUDA kernel in which each block of ``BLOCK_THREADS`` threads
will collectively load, sort, and store its own segment of (``BLOCK_THREADS * ITEMS_PER_THREAD``)
integer keys:
.. code-block:: c++
#include <cub/block/block_load.cuh>
#include <cub/block/block_store.cuh>
#include <cub/block/block_radix_sort.cuh>
template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
__global__ void BlockSortKernel(int *d_in, int *d_out)
{
// Specialize BlockLoad, BlockStore, and BlockRadixSort collective types
using BlockLoadT = cub::BlockLoad<
int, BLOCK_THREADS, ITEMS_PER_THREAD, cub::BLOCK_LOAD_TRANSPOSE>;
using BlockStoreT = cub::BlockStore<
int, BLOCK_THREADS, ITEMS_PER_THREAD, cub::BLOCK_STORE_TRANSPOSE>;
using BlockRadixSortT = cub::BlockRadixSort<
int, BLOCK_THREADS, ITEMS_PER_THREAD>;
// Allocate type-safe, repurposable shared memory for collectives
__shared__ union {
typename BlockLoadT::TempStorage load;
typename BlockStoreT::TempStorage store;
typename BlockRadixSortT::TempStorage sort;
} temp_storage;
// Obtain this block's segment of consecutive keys (blocked across threads)
int thread_keys[ITEMS_PER_THREAD];
const int block_offset = blockIdx.x * (BLOCK_THREADS * ITEMS_PER_THREAD);
BlockLoadT(temp_storage.load).Load(d_in + block_offset, thread_keys);
__syncthreads(); // Barrier for smem reuse
// Collectively sort the keys
BlockRadixSortT(temp_storage.sort).Sort(thread_keys);
__syncthreads(); // Barrier for smem reuse
// Store the sorted segment
BlockStoreT(temp_storage.store).Store(d_out + block_offset, thread_keys);
}
.. code-block:: c++
// Elsewhere in the host program: parameterize and launch a block-sorting
// kernel in which blocks of 128 threads each sort segments of 2048 keys
int *d_in = ...;
int *d_out = ...;
int num_blocks = ...;
BlockSortKernel<128, 16><<<num_blocks, 128>>>(d_in, d_out);
In this example, threads use ``cub::BlockLoad``, ``cub::BlockRadixSort``, and ``cub::BlockStore``
to collectively load, sort and store the block's segment of input items. Because these operations
are cooperative, each primitive requires an allocation of shared memory for threads to communicate
through. The typical usage pattern for a CUB collective is:
#. Statically specialize the primitive for the specific problem setting at hand, e.g.,
the data type being sorted, the number of threads per block, the number of keys per
thread, optional algorithmic alternatives, etc. (CUB primitives are also implicitly
specialized by the targeted compilation architecture.)
#. Allocate (or alias) an instance of the specialized primitive's nested ``TempStorage``
type within a shared memory space.
#. Specify communication details (e.g., the ``TempStorage`` allocation) to
construct an instance of the primitive.
#. Invoke methods on the primitive instance.
In particular, ``cub::BlockRadixSort`` is used to collectively sort the segment of data items
that have been partitioned across the thread block. To provide coalesced accesses
to device memory, we configure the ``cub::BlockLoad`` and ``cub::BlockStore`` primitives
to access memory using a striped access pattern (where consecutive threads
simultaneously access consecutive items) and then *transpose* the keys into
a :ref:`blocked arrangement <flexible-data-arrangement>` of elements across threads.
To reuse shared memory across all three primitives, the thread block statically
allocates a union of their ``TempStorage`` types.
Why do you need CUB?
==================================================
Writing, tuning, and maintaining kernel code is perhaps the most challenging,
time-consuming aspect of CUDA programming. Kernel software is where
the complexity of parallelism is expressed. Programmers must reason about
deadlock, livelock, synchronization, race conditions, shared memory layout,
plurality of state, granularity, throughput, latency, memory bottlenecks, etc.
With the exception of CUB, however, there are few (if any) software libraries of
*reusable* kernel primitives. In the CUDA ecosystem, CUB is unique in this regard.
As a `SIMT <https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#hardware-implementation>`_
library and software abstraction layer, CUB provides:
#. **Simplicity of composition**. CUB enhances programmer productivity by
allowing complex parallel operations to be easily sequenced and nested.
For example, ``cub::BlockRadixSort`` is constructed from ``cub::BlockExchange`` and
``cub::BlockRadixRank``. The latter is composed of ``cub::BlockScan``
which incorporates ``cub::WarpScan``.
.. figure:: ../img/nested_composition.png
:align: center
#. **High performance**. CUB simplifies high-performance program and kernel
development by taking care to implement the state-of-the-art in parallel algorithms.
#. **Performance portability**.
CUB primitives are specialized to match the diversity of NVIDIA hardware, continuously
evolving to accommodate new architecture-specific features and instructions. And
because CUB's device-wide primitives are implemented using flexible block-wide and
warp-wide collectives, we are able to performance-tune them to match the processor
resources provided by each CUDA processor architecture.
#. **Simplicity of performance tuning**:
* **Resource utilization**. CUB primitives allow developers to quickly
change grain sizes (threads per block, items per thread, etc.) to best match
the processor resources of their target architecture
* **Variant tuning**. Most CUB primitives support alternative algorithmic
strategies. For example, ``cub::BlockHistogram`` is parameterized to implement either
an atomic-based approach or a sorting-based approach. (The latter provides uniform
performance regardless of input distribution.)
* **Co-optimization**. When the enclosing kernel
is similarly parameterizable, a tuning configuration can be found that optimally
accommodates their combined register and shared memory pressure.
#. **Robustness and durability**. CUB just works. CUB primitives
are designed to function properly for arbitrary data types and widths of
parallelism (not just for the built-in C++ types or for powers-of-two threads
per block).
#. **Reduced maintenance burden**. CUB provides a SIMT software abstraction layer
over the diversity of CUDA hardware. With CUB, applications can enjoy
performance-portability without intensive and costly rewriting or porting efforts.
#. **A path for language evolution**. CUB primitives are designed
to easily accommodate new features in the CUDA programming model, e.g., thread
subgroups and named barriers, dynamic shared memory allocators, etc.
How do CUB collectives work?
==================================================
Four programming idioms are central to the design of CUB:
#. :ref:`Generic programming <generic-programming>`. C++ templates provide the flexibility
and adaptive code generation needed for CUB primitives to be useful, reusable, and
fast in arbitrary kernel settings.
#. :ref:`Reflective class interfaces <reflective-class-interfaces>`.
CUB collectives statically export their their resource requirements
(e.g., shared memory size and layout) for a given specialization, which allows compile-time
tuning decisions and resource allocation.
#. :ref:`Flexible data arrangement across threads <flexible-data-arrangement>`.
CUB collectives operate on data that is logically partitioned across a group of threads.
For most collective operations, efficiency is increased with increased granularity
(i.e., items per thread).
#. :ref:`Static tuning and co-tuning <static-tuning-and-co-tuning>`. Simple constants and static
types dictate the granularities and algorithmic alternatives to be employed by CUB collectives.
When the enclosing kernel is similarly parameterized, an optimal configuration can be determined
that best accommodates the combined behavior and resource consumption of all primitives within
the kernel.
.. _generic-programming:
Generic programming
--------------------------------------------------
We use template parameters to specialize CUB primitives for the particular
problem setting at hand. Until compile time, CUB primitives are not bound
to any particular:
* Data type (int, float, double, etc.)
* Width of parallelism (threads per thread block)
* Grain size (data items per thread)
* Underlying processor (special instructions, warp size, rules for bank conflicts, etc.)
* Tuning configuration (e.g., latency vs. throughput, algorithm selection, etc.)
.. _reflective-class-interfaces:
Reflective class interfaces
--------------------------------------------------
Unlike traditional function-oriented interfaces, CUB exposes its collective
primitives as templated C++ classes. The resource requirements for a specific
parameterization are reflectively advertised as members of the class. The
resources can then be statically or dynamically allocated, aliased
to global or shared memory, etc. The following illustrates a CUDA kernel
fragment performing a collective prefix sum across the threads of a thread block:
.. code-block:: c++
#include <cub/cub.cuh>
__global__ void SomeKernelFoo(...)
{
// Specialize BlockScan for 128 threads on integer types
using BlockScan = cub::BlockScan<int, 128>;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage scan_storage;
...
// Obtain a segment of consecutive items that are blocked across threads
int thread_data_in[4];
int thread_data_out[4];
...
// Perform an exclusive block-wide prefix sum
BlockScan(scan_storage).ExclusiveSum(thread_data_in, thread_data_out);
Furthermore, the CUB interface is designed to separate parameter
fields by concerns. CUB primitives have three distinct parameter fields:
#. *Static template parameters*. These are constants that will
dictate the storage layout and the unrolling of algorithmic steps (e.g.,
the input data type and the number of block threads), and are used to specialize the class.
#. *Constructor parameters*. These are optional parameters regarding
inter-thread communication (e.g., storage allocation, thread-identifier mapping,
named barriers, etc.), and are orthogonal to the functions exposed by the class.
#. *Formal method parameters*. These are the operational inputs/outputs
for the various functions exposed by the class.
This allows CUB types to easily accommodate new
programming model features (e.g., named barriers, memory allocators, etc.)
without incurring a combinatorial growth of interface methods.
.. _flexible-data-arrangement:
Flexible data arrangement across threads
--------------------------------------------------
CUDA kernels are often designed such that each thread block is assigned a
segment of data items for processing.
.. figure:: ../img/tile.png
:align: center
:alt: Segment of eight ordered data items
:name: fig_tile
Segment of eight ordered data items
When the tile size equals the thread block size, the
mapping of data onto threads is straightforward (one datum per thread).
However, there are often performance advantages for processing more
than one datum per thread. Increased granularity corresponds to
decreased communication overhead. For these scenarios, CUB primitives
will specify which of the following partitioning alternatives they
accommodate:
.. list-table::
:class: table-no-stripes
:widths: 70 30
* - **Blocked arrangement**. The aggregate tile of items is partitioned
evenly across threads in "blocked" fashion with *thread*\ :sub:`i`
owning the *i*\ :sup:`th` segment of consecutive elements.
Blocked arrangements are often desirable for algorithmic benefits (where
long sequences of items can be processed sequentially within each thread).
- .. figure:: ../img/blocked.png
:align: center
:alt: *Blocked* arrangement across four threads
:name: fig_blocked
*Blocked* arrangement across four threads
(emphasis on items owned by *thread*\ :sub:`0`)
* - **Striped arrangement**. The aggregate tile of items is partitioned across threads in "striped"
fashion, i.e., the ``ITEMS_PER_THREAD`` items owned by each thread have logical stride
``BLOCK_THREADS`` between them. Striped arrangements are often desirable for data movement through
global memory (where
`read/write coalescing <https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#coalesced-access-to-global-memory>`_
is an important performance consideration).
- .. figure:: ../img/striped.png
:align: center
:alt: *Striped* arrangement across four threads
:name: fig_striped
*Striped* arrangement across four threads
(emphasis on items owned by *thread*\ :sub:`0`)
The benefits of processing multiple items per thread (a.k.a., *register blocking*,
*granularity coarsening*, etc.) include:
* Algorithmic efficiency. Sequential work over multiple items in
thread-private registers is cheaper than synchronized, cooperative
work through shared memory spaces.
* Data occupancy. The number of items that can be resident on-chip in
thread-private register storage is often greater than the number of
schedulable threads.
* Instruction-level parallelism. Multiple items per thread also
facilitates greater ILP for improved throughput and utilization.
Finally, ``cub::BlockExchange`` provides operations for converting between blocked
and striped arrangements.
.. _static-tuning-and-co-tuning:
Static tuning and co-tuning
--------------------------------------------------
This style of flexible interface simplifies performance tuning. Most CUB
primitives support alternative algorithmic strategies that can be
statically targeted by a compiler-based or JIT-based autotuner. (For
example, ``cub::BlockHistogram`` is parameterized to implement either an
atomic-based approach or a sorting-based approach.) Algorithms are also
tunable over parameters such as thread count and grain size as well.
Taken together, each of the CUB algorithms provides a fairly rich tuning
space.
Whereas conventional libraries are optimized offline and in isolation, CUB
provides interesting opportunities for whole-program optimization. For
example, each CUB primitive is typically parameterized by threads-per-block
and items-per-thread, both of which affect the underlying algorithm's
efficiency and resource requirements. When the enclosing kernel is similarly
parameterized, the coupled CUB primitives adjust accordingly. This enables
autotuners to search for a single configuration that maximizes the performance
of the entire kernel for a given set of hardware resources.
How do I get started using CUB?
==================================================
CUB is a C++ header-only library, and part of the CUDA Core Compute Libraries (CCCL).
It ships as part of the CUDA Toolkit and is thus readily available when using the ``nvcc`` compiler.
Alternatively, consider fetching CCCL directly from GitHub to benefit from the latest improvements.
How is CUB different than Thrust and Modern GPU?
==================================================
CUB and Thrust
--------------------------------------------------
CUB and :ref:`Thrust <thrust-module>` share some
similarities in that they both provide similar device-wide primitives for CUDA.
However, they target different abstraction layers for parallel computing.
Thrust abstractions are agnostic of any particular parallel framework (e.g.,
CUDA, TBB, OpenMP, sequential CPU, etc.). While Thrust has a "backend"
for CUDA devices, Thrust interfaces themselves are not CUDA-specific and
do not explicitly expose CUDA-specific details (e.g., ``cudaStream_t`` parameters).
CUB, on the other hand, is slightly lower-level than Thrust. CUB is specific
to CUDA C++ and its interfaces explicitly accommodate CUDA-specific features.
Furthermore, CUB is also a library of SIMT collective primitives for block-wide
and warp-wide kernel programming.
CUB and Thrust are complementary and can be used together. In fact, the CUB
project arose out of a maintenance need to achieve better performance-portability
within Thrust by using reusable block-wide primitives to reduce maintenance and
tuning effort.
CUB and Modern GPU
--------------------------------------------------
CUB and `Modern GPU <https://github.com/moderngpu/moderngpu>`_ also
share some similarities in that they both implement similar device-wide primitives for CUDA.
However, they serve different purposes for the CUDA programming community. MGPU
is a pedagogical tool for high-performance GPU computing, providing clear and concise
exemplary code and accompanying commentary. It serves as an excellent source of
educational, tutorial, CUDA-by-example material. The MGPU source code is intended
to be read and studied, and often favors simplicity at the expense of portability and
flexibility.
CUB, on the other hand, is a production-quality library whose sources are complicated
by support for every version of CUDA architecture, and is validated by an extensive
suite of regression tests. Although well-documented, the CUB source text is verbose
and relies heavily on C++ template metaprogramming for situational specialization.
CUB and MGPU are complementary in that MGPU serves as an excellent descriptive source
for many of the algorithmic techniques used by CUB.
Contributors
==================================================
CUB is developed as open-source as part of the CUDA Core Compute Libraries (CCCL) by NVIDIA.
Open Source License
==================================================
CUB is mostly licensed under the BSD 3-Clause "New" or "Revised" License.
New files are created under the Apache-2.0 WITH LLVM-exception License.
See also our `LICENSE <https://github.com/NVIDIA/cccl/blob/main/LICENSE>`_ file.

8
cccl_upstream/docs/cub/thread_level.rst Normal file
View File

@@ -0,0 +1,8 @@
.. _thread-module:
Thread-level Primitives
==================================================
CUB thread-level algorithms are specialized for execution by a single thread.
* :cpp:func:`cub::ThreadReduce <cub::ThreadReduce>` computes reduction of a sequence of items

146
cccl_upstream/docs/cub/tuning.rst Normal file
View File

@@ -0,0 +1,146 @@
..
TODO(bgruber): rename the label below to _cub-tuning when all tuning API exposure PRs have landed
.. _cub-policy-selectors:
Tunings
================================================================================
Device-scope algorithms in CUB have many knobs that significantly impact performance (without affecting correctness).
For instance, the number of threads per block and items per thread can be tuned to maximize performance for a given device and data type.
But also algorithmic choices such as the load or store algorithm, a load vectorization size,
the used block-level algorithm, or enabling the use of the tensor memory accelerator can be tweaked.
Most device-scope algorithms in CUB accept a set of such tuning parameters,
by passing a policy selector, wrapped into :code:`cuda::execution::tune(...)`,
as part of the environment of a CUB device-scope API.
In the following, we describe this process in more detail.
Policy selectors
--------------------------------------------------------------------------------
Policy selectors are the mechanism by which CUB's device algorithms and kernels select
tuning parameters for a given workload and GPU compute capability.
A policy selector is a stateless callable that maps a :code:`cuda::compute_capability` to a policy struct
containing the tuning values for that compute capability.
Each set of CUB algorithms using a common underlying implementation defines a common policy struct,
e.g. :code:`cub::ReducePolicy`, which must be returned by a policy selector passed to those algorithms.
CUB employs internal default policy selectors providing tunings for known compute capabilities and workloads,
which are not publicly accessible to users.
Users can override CUB's policy selector for a given algorithm
by passing a custom policy selector through the algorithm's environment parameter.
For a description of how policy selectors are used internally in CUB's dispatch layer and kernels
see the corresponding :ref:`developer documentation <cub-developer-guide-device-scope>`.
Defining a policy selector
--------------------------------------------------------------------------------
A policy selector is any type with a :code:`__host__`, :code:`__device__`, :code:`constexpr`, and :code:`const` call operator
taking a ``cuda::compute_capability`` and returning the algorithm's policy struct:
.. code:: c++
template <typename T>
struct my_reduce_tuning {
__host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::ReducePolicy {
// tuning for Hopper and later
if (cc >= cuda::compute_capability(9, 0)) {
const auto pass = cub::ReducePassPolicy{
.threads_per_block = 512,
.items_per_thread = std::max(64 / sizeof(T), 1), // 8 double, 16 float, 32 half_t, ...
.vec_size = 2,
.reduce_algorithm = cub::BLOCK_REDUCE_WARP_REDUCTIONS,
.load_modifier = cub::LOAD_DEFAULT
};
return { .multi_tile = pass, .single_tile = pass };
}
// fallback for older GPUs
const auto pass = cub::ReducePassPolicy{
.threads_per_block = 256,
.items_per_thread = 12,
.vec_size = 1,
.reduce_algorithm = cub::BLOCK_REDUCE_WARP_REDUCTIONS,
.load_modifier = cub::LOAD_DEFAULT
};
return { .multi_tile = pass, .single_tile = pass };
}
};
.. warning::
The policy selector must be stateless (:code:`std::is_empty_v<T>` must be :code:`true`)
since only its type will be passed to a kernel — any captured state would be silently
lost. Policy selectors are also freely constructed and copied where needed, so they
must also be default constructible and copyable (:code:`std::semiregular<T>` must be
:code:`true`).
It can be a class template, but then only a full specialization can be passed to CUB,
e.g., :code:`my_reduce_tuning<float>`.
The implementation can use branches and helper functions arbitrarily,
as long as they can be evaluated at compile-time.
The returned tuning policy can, and probably should, contain different values for different compute capabilities or workloads.
Each CUB algorithm with an environment parameter searches the environment for a policy selector returning a matching policy struct.
If one is found, the policy selector will be used to determine the tuning values for host-side dispatch and kernel compilation.
Each CUB algorithm documents the policy struct to which it responds.
If CUB does not find a matching policy selector in the environment, it falls back to its internal default policy selector.
Multiple policy selectors returning different policy structs can be passed as part of the same environment,
and each algorithm will pick the one with the matching policy struct.
This is useful if the same environment is reused across several algorithm calls.
The policy structs themselves are simple semiregular aggregates.
They support C++20 designated initializers (i.e., the syntax :code:`{ .threads_per_block = 512, ... }`),
comparison for (in-)equality, and serialization using :code:`operator<<`.
They may occasionally contain member functions that compute derived values from the contained tuning values.
All policy structs are public types and will evolve in a non-breaking way, at least during minor releases,
by only adding new data members at the end of the struct.
Passing a policy selector to CUB device-scope algorithms
--------------------------------------------------------------------------------
Custom policy selectors are passed to CUB algorithms via the environment argument.
They first need to be wrapped by passing them to :code:`cuda::execution::tune(...)`:
.. code:: c++
cub::DeviceReduce::Reduce(
d_in, d_out, num_items, op, init,
cuda::execution::tune(my_reduce_tuning<int>{}));
Multiple tunings for different algorithms can be combined in a single environment:
.. code:: c++
auto env = cuda::execution::tune(
my_reduce_tuning<int>{},
my_scan_tuning{}
);
cub::DeviceReduce::Reduce(d_in, d_out, num_items, op, init, env);
cub::DeviceScan::ExclusiveSum(d_in, d_out, num_items, env);
An environment can carry further properties like a stream or a memory resource.
Policy selectors can simply be added to those:
.. code:: c++
auto env = cuda::std::execution::env{
stream_ref,
resource,
cuda::execution::tune(
my_reduce_tuning<int>{}, my_scan_tuning{})
};
cub::DeviceReduce::Reduce(d_in, d_out, num_items, op, init, env);
// alternatively, if we want to extend an env `other_env`
auto env = cuda::std::execution::env{
other_env,
cuda::execution::tune(my_scan_tuning{})
};
cub::DeviceReduce::Reduce(d_in, d_out, num_items, op, init, env);
CUB's benchmarks also make heavy use of policy selectors for tuning.
For more details on authoring benchmarks and their policy selectors for automatic tuning, see :ref:`cub-tuning-infra`.

629
cccl_upstream/docs/cub/tuning_infra.rst Normal file
View File

@@ -0,0 +1,629 @@
.. _cub-tuning-infra:
Automated Tuning Infrastructure
================================================================================
This page describes tuning infrastructure to automatically tune CUB device-scope algorithms for performance.
It provides a set of tools facilitating the process of selecting optimal tuning parameters for a given device and data type.
Terminology
--------------------------------------------------------------------------------
*We omit the word "tuning" but assume it in the definitions for all terms below,
so those terms may mean something else in a more generic context.*
The following three terms are fundamental to understanding the essence of CUB tuning:
* **compile-time (ct) workload**: a workload that can be recognized only at compile time.
*e.g. the combination of key type and offset type,* :code:`int16_t` *and* :code:`int32_t`
* **runtime (rt) workload**: a workload that can be recognized only at runtime.
*e.g. the number of input elements*
* **tuning parameter (or parameter)**: a parameter that can be tuned to maximize performance for a given device and data type.
*e.g. number of threads per block, items per thread*
Algorithms are tuned for different workloads. These workloads are defined as subspaces by NVBench via the benchmarks' axis.
For instance, radix sort can be tuned for different key types, different number of keys, and different distributions of keys. The tuning process is summarized in
the following statement:
.. raw:: html
<div style="display: flex; justify-content: center; align-items: center; height: 10; text-align: center; font-size: 1.5em; color: #76B900;">
"For each Compile-time Workload, we search for the best tuning parameters"
</div>
More specifically, the tuning infrastructure optimizes algorithms for specific compile-time workloads,
aggregating results across all runtime workloads.
It searches through a space of parameters to find the combination for a given compile-time workload with the highest score.
--------
Following is supplemental terminology that will be used throughout the rest of this tuning guide:
.. e.g. :math:`threads\_per\_block=128`
* **Parameter Space**: the set of all possible values for a given Tuning Parameter.
*It is specific to the algorithm*. For example the parameter space for the number of threads per block can be :math:`\{32, 64, 96, 128, \dots, 1024\}` for radix sort, but :math:`\{32, 64, 128, 256, 512\}` for merge sort.
* **Search Space**: Cartesian product of all the Parameter Spaces of a single algorithm.
For instance, the Search Space for an algorithm with tunable items per thread and threads per block might look like :math:`\{(ipt \times tpb) | ipt \in \{1, \dots, 25\} \text{and} tpb \in \{32, 64, 96, 128, \dots, 1024\}\}`.
* **Variant** - a point in the corresponding Search Space.
* **Base** - the variant that CUB uses by default.
* **Score** - a single number representing the performance for a given compile-time workload across all runtime workloads. For instance, a weighted-sum of speedups of a given variant compared to its base for all runtime workloads is a score.
.. * **Search** - a process consisting of covering all variants for all compile-time workloads to find a variant with maximal score.
.. ^^^ @giannis: again we do not want to scare a first time user with too many terms. "search" is both evident and can also be explained with an introductory sentence in the "Search Process" chapter ^^^
.. _cub-tuning-infra-authoring-benchmarks:
Authoring Benchmarks
--------------------------------------------------------------------------------
CUB benchmarks are split into multiple files based on the algorithm they are testing
and potentially further into compile-time flavors that are tuned for individually
(e.g. sorting only keys vs. key-value pairs, or reducing using sum vs. using min).
The name of the directory represents the name of the algorithm.
The filename corresponds on the flavor.
For instance, the benchmark :code:`benchmarks/bench/radix_sort/keys.cu` tests the radix sort implementation sorting only keys.
The executable file name is going to be transformed into :code:`cub.bench.radix_sort.keys.*`,
which is the benchmark name reported by the infrastructure.
+++++++++++++++
Headers
+++++++++++++++
**Benchmarks are based on NVBench.**
You start writing a benchmark by including :code:`nvbench_helper.cuh`. This contains all
necessary includes and definitions.
.. code:: c++
#include <nvbench_helper.cuh>
The next step is to define a search space. The search space is represented by a number of C++ comments.
The format consists of the :code:`%RANGE%` keyword, the parameter macro, the parameter abbreviation, and its range of values.
The range is represented by three numbers: :code:`start:end:step`.
Start and end are included.
For instance, the following code defines a search space for two parameters, the number of threads per block and items per thread.
.. code:: c++
// %RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1
// %RANGE% TUNE_THREADS_PER_BLOCK tpb 128:1024:32
Next, you need to define a benchmark function. The function accepts :code:`nvbench::state &state` and
a :code:`nvbench::type_list`. For more details on the benchmark signature, take a look at the
`NVBench documentation <https://github.com/NVIDIA/nvbench>`_.
.. code:: c++
template <typename T, typename OffsetT>
void algname(nvbench::state &state, nvbench::type_list<T, OffsetT>)
{...}
Before proceeding further with the benchmark authoring
it is imperative to understand policy selectors and how they provide tuning values.
+++++++++++++++
Policy Selector
+++++++++++++++
The tuning value of all CUB device algorithms can be customized
by providing a custom policy selector to the environment argument of a CUB API call.
See :ref:`cub-policy-selectors` for a full explanation.
The tuning infrastructure will use the :code:`TUNE_BASE` macro to distinguish between compiling the base version (i.e. baseline) of a benchmark
and compiling a variant for a given set of tuning parameters.
When base is used, no custom policy selector is specified, so CUB's default tunings are used.
If :code:`TUNE_BASE` is not defined, we define a custom policy selector
that specifies the values for the current variant (i.e. the current set of tuning parameters),
which are derived from the parameter macros defined in the :code:`%RANGE%` comments, which define the search space.
This custom policy selector is passed to :code:`cuda::execution::tune`,
and the returned value is included in the environment passed to the CUB API.
The following code is included in the benchmark for the policy selector to be enabled
and the parameters to have effect in execution:
.. code:: c++
#if !TUNE_BASE
template <typename T>
struct policy_selector {
_CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability /*cc*/) const -> cub::ReducePolicy {
const auto [items, threads] = cub::detail::scale_mem_bound(
TUNE_THREADS_PER_BLOCK, TUNE_ITEMS_PER_THREAD, sizeof(T));
const auto pass = cub::ReducePassPolicy{
.threads_per_block = threads,
.items_per_thread = items,
.vec_size = 1 << TUNE_ITEMS_PER_VEC_LOAD_POW2,
.reduce_algorithm = cub::BLOCK_REDUCE_WARP_REDUCTIONS,
.load_modifier = cub::LOAD_DEFAULT
};
return { .multi_tile = pass, .single_tile = pass};
}
};
#endif
The custom policy selector returns a fixed policy regardless of the compute capability,
since tuning is usually done on a specific GPU and compiling only for that GPU's compute capability.
The policy selector uses all tuning parameters from the search space to form the policy used by CUB.
+++++++++
Main Body
+++++++++
The :code:`state` passed into the benchmark function allows access to runtime workload axes,
for example the number of elements to process.
*When creating containers for the input avoid to initialize data yourself.
Instead, use the* :code:`gen` *function,
which will fill the input vector with random data on GPU with no compile-time overhead.*
.. code:: c++
const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}"));
thrust::device_vector<T> in(elements);
thrust::device_vector<T> out(1);
gen(seed_t{}, in);
In addition to the benchmark runtime, NVBench can also report information on the achieved memory bandwidth.
For this, you can optionally provide information on the memory reads and writes of the algorithm to the :code:`state`:
.. code:: c++
state.add_element_count(elements);
state.add_global_memory_reads<T>(elements, "Size");
state.add_global_memory_writes<T>(1);
Finally, the actual benchmark region then calls the public CUB API with an appropriate environment.
Temporary storage allocation is handled automatically by a caching allocator provided through the environment.
The CUDA stream is provided by the :code:`nvbench::launch` parameter of the benchmark lambda.
When tuning (:code:`TUNE_BASE` is not defined),
an instance of the custom policy selector is wrapped by :code:`cuda::execution::tune` and passed to the environment as well.
.. code:: c++
caching_allocator_t alloc;
state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch,
[&](nvbench::launch &launch) {
auto env = cub_bench_env(
alloc,
launch
#if !TUNE_BASE
, cuda::execution::tune(policy_selector<T>{})
#endif // !TUNE_BASE
);
cub::DeviceReduce::Reduce(d_in, d_out, elements, op_t{}, init, env);
});
This concludes defining the benchmark function.
Now we need to tell NVBench about it.
++++++++++++++++++
NVBench Attributes
++++++++++++++++++
.. code:: c++
NVBENCH_BENCH_TYPES(algname, NVBENCH_TYPE_AXES(all_types, offset_types))
.set_name("base")
.set_type_axes_names({"T{ct}", "OffsetT{ct}"})
.add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4));
:code:`NVBENCH_BENCH_TYPES` registers the benchmark as one with multiple compile-time workloads,
which are defined by the Cartesian product of the type lists in :code:`NVBENCH_TYPE_AXES`.
:code:`set_name(...)` sets the name of the benchmark.
Only alphabetical characters, numbers and underscores are allowed in the benchmark name.
Furthermore, compile-time axes should be suffixed with :code:`{ct}`. The runtime axes might be optionally annotated
as :code:`{io}` which stands for importance-ordered. *This will tell the tuning infrastructure that
the later values on the axis are more important. If the axis is not annotated, each value will be
treated as equally important.*
When you define a type axis annotated with :code:`{ct}`, you should consider optimizing
the build time. Many variants are going to be build, but the search is considering one compile-time
use case at a time. This means that if you have many types to tune for, you'll end up having
many template specializations that you don't need. To avoid this, for each compile time axis, the tuning framework will predefine
a `TUNE_AxisName` macro with the type that's currently being tuned. For instance, if you
have the type axes :code:`T{ct}` and :code:`OffsetT` (as shown above), you can use the following
pattern to narrow down the types you compile for:
.. code:: c++
#ifdef TUNE_T
using all_types = nvbench::type_list<TUNE_T>;
#else
using all_types = nvbench::type_list<char, short, int, long, ...>;
#endif
#ifdef TUNE_OffsetT
using offset_types = nvbench::type_list<TUNE_OffsetT>;
#else
using offset_types = nvbench::type_list<int32_t, int64_t>;
#endif
This logic is already implemented if you use any of the following predefined type lists:
.. list-table:: Predefined type lists
:header-rows: 1
* - Axis name
- C++ identifier
- Included types
* - :code:`T{ct}`
- :code:`integral_types`
- :code:`int8_t, int16_t, int32_t, int64_t`
* - :code:`T{ct}`
- :code:`fundamental_types`
- :code:`integral_types` and :code:`int128_t, float, double`
* - :code:`T{ct}`
- :code:`all_types`
- :code:`fundamental_types` and :code:`complex`
* - :code:`OffsetT{ct}`
- :code:`offset_types`
- :code:`int32_t, int64_t`
You are free to define your own axis names and use the logic above for them (see the sort pairs example).
A single benchmark file can define multiple benchmarks (multiple benchmark functions registered with :code:`NVBENCH_BENCH_TYPES`).
All benchmarks in a single file must share the same compile-time axes.
**The tuning infrastructure will run all benchmarks in a single file together for the same compile-time workload
and compute a common score across all benchmarks and runtime workloads.
Unless a benchmark axis is importance-ordered, each sample contributes equally to the score.**
This is useful to tune an algorithm for multiple runtime use cases at once,
that we don't intend to provide separate tuning policies for.
Also, a large space of runtime workloads can be segmented this way,
e.g. by splitting the benchmark entry point and supplying a few low and a few high values for a runtime axis:
.. code:: c++
NVBENCH_BENCH_TYPES(algname, NVBENCH_TYPE_AXES(all_types, offset_types))
.set_name("small")
...
.add_int64_power_of_two_axis("SegmentSize", nvbench::range(0, 3, 1)); // tests sizes 2^0, 2^1, 2^2, 2^3
NVBENCH_BENCH_TYPES(algname, NVBENCH_TYPE_AXES(all_types, offset_types))
.set_name("large")
...
.add_int64_power_of_two_axis("SegmentSize", nvbench::range(12, 18, 2)); // tests sizes 2^12, 2^14, 2^16, 2^18
Search Process
--------------------------------------------------------------------------------
During the Search Process we are covering all variants for all compile-time workloads to find a variant with a maximum (at least locally) score.
To get started with tuning, you need to configure CMake.
You can use the following command:
.. code:: bash
$ mkdir build
$ cd build
$ cmake .. --preset=cub-tune
You can then run the tuning search for a specific algorithm and compile-time workload. We use a CCCL internal script for that:
.. code:: bash
$ ../benchmarks/scripts/search.py -R '.*merge_sort.*pairs' -a 'KeyT{ct}=I128' -a 'Elements{io}[pow2]=28'
cub.bench.merge_sort.pairs.trp_0.ld_1.ipt_13.tpb_6 0.6805093269929858
cub.bench.merge_sort.pairs.trp_0.ld_1.ipt_11.tpb_10 1.0774560502969677
...
This will search the space of merge sort for key-value pairs, for the key type :code:`int128_t` on :code:`2^28` elements.
The :code:`-R` and :code:`-a` options are optional. **If not specified, all benchmarks are going to be tuned.**
The :code:`-R` option can select multiple benchmarks using a regular expression.
For the axis option :code:`-a`, you can also specify a range of values like :code:`-a 'KeyT{ct}=[I32,I64]'`.
Any axis values not supported by a selected benchmark will be ignored.
The first variant :code:`cub.bench.merge_sort.pairs.trp_0.ld_1.ipt_13.tpb_6` has a score <1 and is thus generally slower than the baseline,
whereas the second variant :code:`cub.bench.merge_sort.pairs.trp_0.ld_1.ipt_11.tpb_10` has a score of >1 and is thus an improvement over the baseline.
.. warning::
Notice there is currently a limitation in :code:`search.py`
which will only execute runs for the first axis value for each axis
(independently of whether the axis is specified on the command line or not).
Tuning for multiple axis values requires multiple runs of :code:`search.py`.
Please see `this issue <https://github.com/NVIDIA/cccl/issues/2267>`_ for more information.
**Benchmarks do not need to be built a priori.** The tuning framework will handle building the benchmarks (base and variants) and running them by itself.
It will keep track of the build time for base and variants.
Sometimes, a tuning variant may lead the compiler to hang or take exceptionally long to compile.
To keep the tuning process going, if the build time of a variant exceeds a threshold, the build is cancelled.
The same applies to benchmarks running for too long.
To get quick feedback on what benchmarks are selected and how big the search space is,
you can add the :code:`-l` option:
.. code:: bash
$ ../benchmarks/scripts/search.py -R '.*merge_sort.*pairs' -a 'KeyT{ct}=I128' -a 'Elements{io}[pow2]=28' -l
ctk: 12.6.85
cccl: v2.7.0
### Benchmarks
* `cub.bench.merge_sort.pairs`: 540 variants:
* `trp`: (0, 2, 1)
* `ld`: (0, 3, 1)
* `ipt`: (7, 25, 1)
* `tpb`: (6, 11, 1)
It will list all selected benchmarks as well as the total number of variants (the magnitude of the search space)
as a result of the Cartesian product of all its tuning parameter spaces.
The tuning infrastructure stores the results in an SQLite database called :code:`cccl_meta_bench.db` in the build directory.
This database persists across tuning runs.
If you interrupt the benchmark script and then launch it again, only missing benchmark variants will be run.
Tuning on multiple GPUs
--------------------------------------------------------------------------------
Because the search process computes scores by comparing the performance of a variant to the baseline,
it has to store the baseline result in the tuning database.
The baseline is specific to the physical GPU on which it was obtained.
Therefore, a single tuning database should not be used to run the tuning search on two different GPUs, even of the same architecture.
Similarly, you should also not interrupt the search and resume it on a different GPU.
Be careful when sharing build directories over network file systems.
Check whether a build directory already contains a :code:`cccl_meta_bench.db` from a previous run before starting a new search.
..
TODO(bgruber): I don't yet understand whether we can tune a single variant on multiple GPUs.
I think this is possible, but would it then create a database per GPU (because 1 baseline per GPU)?
Does search.py do this automatically, or do I need to pass a flag? Or does this only work with our "internal extensions"?
Because the search space can be separated based on different axis values,
a tuning search can be run on multiple GPUs in parallel, even across multiple physical machines (e.g., on a cluster).
To do this, :code:`search.py` is invoked in parallel, one invocation/process per GPU,
with different axis values specified for each invocation.
A dedicated tuning database will be created per physical GPU.
If a shared filesystem is in use, make sure that :code:`search.py` is run from different directories,
so the :code:`cccl_meta_bench.db` files are placed into distinct paths.
It is recommended to drive a multi-GPU/multi-node search process from a script,
iterating the axis values and invoking :code:`search.py` for each variant.
This integrates nicely with workload managers on clusters, which allow submitting batch jobs.
In such a scenario, it is recommended to submit a job per variant.
After tuning on multiple GPUs, the results are available in multiple tuning databases, which can be analyzed together.
Analyzing the results
--------------------------------------------------------------------------------
The result of the search is stored in one or more :code:`cccl_meta_bench.db` files. To analyze the
result you can use the :code:`analyze.py` script.
The :code:`--coverage` flag will show the amount of variants that were covered per compile-time workload:
.. code:: bash
$ ../benchmarks/scripts/analyze.py --coverage
cub.bench.radix_sort.keys[T{ct}=I8, OffsetT{ct}=I32] coverage: 167 / 522 (31.9923%)
cub.bench.radix_sort.keys[T{ct}=I8, OffsetT{ct}=I64] coverage: 152 / 522 (29.1188%)
The :code:`--top N` flag will list the best :code:`N` variants for each compile-time workload:
.. code:: bash
$ ../benchmarks/scripts/analyze.py --top=5
cub.bench.radix_sort.keys[T{ct}=I8, OffsetT{ct}=I32]:
variant score mins means maxs
97 ipt_19.tpb_512 1.141015 1.039052 1.243448 1.679558
84 ipt_18.tpb_512 1.136463 1.030434 1.245825 1.668038
68 ipt_17.tpb_512 1.132696 1.020470 1.250665 1.688889
41 ipt_15.tpb_576 1.124077 1.011560 1.245011 1.722379
52 ipt_16.tpb_512 1.121044 0.995238 1.252378 1.717514
cub.bench.radix_sort.keys[T{ct}=I8, OffsetT{ct}=I64]:
variant score mins means maxs
71 ipt_19.tpb_512 1.250941 1.155738 1.321665 1.647868
86 ipt_20.tpb_512 1.250840 1.128940 1.308591 1.612382
55 ipt_17.tpb_512 1.244399 1.152033 1.327424 1.692091
98 ipt_21.tpb_448 1.231045 1.152798 1.298332 1.621110
85 ipt_20.tpb_480 1.229382 1.135447 1.294937 1.631225
The name of the variant contains the short parameter names and values used for the variant.
For each variant, a score is reported. The base has a score of 1.0, so each score higher than 1.0 is an improvement over the base.
However, because a single variant contains multiple runtime workloads, also the minimum, mean, maximum score is reported.
If all those three values are larger than 1.0, the variant is strictly better than the base.
If only the mean or max are larger than 1.0, the variant may perform better in most runtime workloads, but regress in others.
This information can be used to change the existing tuning policies in CUB. A detailed explanation of the output is presented
in the following image:
.. image:: ../images/top_results_expl.png
By default, :code:`analyze.py` will look for a file named :code:`cccl_meta_bench.db` in the current directory.
If the tuning results are available in multiple databases, e.g., after tuning on multiple GPUs,
glob expressions matching multiple databases, or just multiple file paths, can be passed as arguments as well:
.. code:: bash
$ ../benchmarks/scripts/analyze.py --top=5 <path-to-databases>/*.db
In case the tuning database(s) store(s) results for several different benchmarks,
the analysis can again be restricted using a regular expression via the :code:`-R` option:
.. code:: bash
$ ../benchmarks/scripts/analyze.py -R=".*radix_sort.keys.*" --top=5 <path-to-databases>/*.db
Variant plots
--------------------------------------------------------------------------------
The reported score for a tuning aggregates the performance across all runtime workloads.
Furthermore, NVBench collects and aggregates multiple samples for a single compile and runtime workload.
So, even though the min, mean and max score are reported for a variant,
it may be necessary to compare the distributions of raw speedups between the baseline and a variant across all runtime workloads and samples.
This is achieved using variant plots.
For more background information on this subject, we refer the reader to `this article <https://aakinshin.net/posts/shift-and-ratio-functions/>`_.
A variant plot can be generated for one or more variants using the :code:`--variants-ratio=` option and specifying the specific variant to plot.
For example:
.. code:: bash
$ ../benchmarks/scripts/analyze.py -R=".*radix_sort.keys.*" --variants-ratio='ipt_18.tpb_288' <path-to-databases>/*.db
May display a matrix of variant plots like:
.. image:: ../images/variant_plot.png
In the image above we see twelve diagrams for the Cartesian product of the :code:`Entropy` (horizontally) and :code:`Elements{io}` (vertically) runtime axes.
The compile-time axes are fixed for one matrix of variant plots.
Across each variant plot's x-axis, the speedup over the baseline (y-axis) is represented.
The baseline is shown as a straight horizontal red line at 1.
The found tuning thus results in a slowdown for :code:`Elements{io}` 2^16 and 2^20 (orange line below red baseline),
but a speedup for 2^24 and 2^28 (orange line above red baseline).
In general, bigger axis values for plots for importance-ordered axes, like :code:`Elements{io}`,
should be prioritized in evaluating a given tuning, because GPUs are optimized for large problem sizes.
However, while the almost 4% slowdown for 2^16 elements at entropy 0.544 may be bearable,
a close to 7% slowdown for 2^20 elements at entropy 1 is probably too large to accept this tuning,
despite the solid 3.5-8% speedup for larger element counts.
The shown ratios are generated by fitting an equal amount of quantiles into the samples of the baseline and the variant,
and then showing the quotient for each corresponding quantile from baseline and variant.
For background information on the quantile-respectful density estimation,
we refer the reader to this `article <https://aakinshin.net/posts/qrde-hd>`_.
By default, a quantile corresponds to a percentile, and thus a ratio plot contains 100 data points
expressing the speedup of the slowest 1% in the variant over the slowest 1% in the baseline (left),
then the second slowest 1%, etc., until the speedup of the fastest 1% in the variant over the fastest 1% in the baseline (right).
The detailed analysis via variant plots is needed,
because a single aggregated score cannot represent the distribution of samples obtained from highly concurrent algorithms, such as those in CUB.
Even though NVBench reruns a benchmark many times to gain statistical confidence in the result,
the runtime of a CUB algorithm does not necessarily follow a normal distribution.
For example, the concurrent nature of some algorithms may result in bimodal or even more complex distributions,
as a consequence of how the hardware schedules and executes threads.
Also, the kind of distribution may be different between baseline and variant.
For all these reasons, comparing the distribution of samples is the only reliable way to determine,
whether a tuning provides a consistent speedup for all runtime workloads.
Creating and extending tuning policies
--------------------------------------------------------------------------------
Once a suitable tuning result has been selected, we have to translate it into C++ code that will be picked up by CUB.
The tuning variant name shown by :code:`analyze.py` gives us all the information on the selected tuning values.
Here is an example:
.. code:: bash
$ ../benchmarks/scripts/analyze.py --top=1
cub.bench.radix_sort.keys[T{ct}=I8, OffsetT{ct}=I64]:
variant score mins means maxs
71 ipt_19.tpb_512 1.250941 1.155738 1.321665 1.647868
Assume we have determined this tuning to be the best one for sorting I8 keys using radix_sort using I64 offsets.
The ``variant`` can be decoded using the ``// %RANGE%`` comments in the C++ source code of the benchmark,
since the names of the reported parameters in the variant are derived from these:
.. code:: c++
// %RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1
// %RANGE% TUNE_THREADS_PER_BLOCK tpb 128:1024:32
The variant ``ipt_19.tpb_512``, which stands for 19 items per thread (``ipt``) and 512 threads per block (``tpb``),
was thus compiled with ``-DTUNE_ITEMS_PER_THREAD=19 -DTUNE_THREADS_PER_BLOCK=512``.
The meaning of these values is specific to the benchmark definition,
and we have to check the benchmark's source code for how they are applied.
Equally named tuning parameters may not translate to different benchmarks (please double check).
As a user of CUB, such a new set of tuning parameters (i.e. a variant) can then be used to define a policy selector,
which is passed to the public CUB API through the environment,
as :ref:`sketched above <cub-tuning-infra-authoring-benchmarks>`:
.. code:: c++
struct policy_selector {
_CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability /*cc*/) const -> cub::AlgorithmPolicy {
return {
.threads_per_block = 512,
.items_per_thread = 19,
...
};
}
};
The default tunings defined inside CUB's source use the same infrastructure
but should only be changed and extended by the CCCL maintainers.
All default tunings are found in the :code:`cub/device/dispatch/tuning/tuning_*.cuh` headers, organized by algorithm.
CUB's policy selectors are highly parameterized on type information and traits of the input arguments to CUB algorithms
(like accumulator type, offset size, and operation kind),
which they turn into a policy for a given compute capability.
The way tuning values are selected is different for each CUB algorithm and requires studying the corresponding code.
The general principles of policy selectors and tunings are documented :ref:`here <cub-policy-selectors>`.
For example, signed and unsigned integers of the same size are often represented by the same tuning.
In general, variants for which the algorithmic behavior is expected to be the same
(same arithmetic intensity, no special instructions for one of the data types, same amount of bytes to load/store, etc.)
are covered by the same tuning.
When a better variant has been found and CUB already has a tuning for this variant,
the tuning parameter values can simply be updated in the corresponding CUB tuning header.
This is usually the case when a CUB algorithm has been reengineered and shows different performance characteristics,
or more tuning parameters are exposed (e.g., a new load algorithm is available).
For example, an existing tuning selection function may contain code like:
.. code:: c++
constexpr auto get_sm90_tuning(type_t accum_t, op_kind_t op, int offset_size, int accum_size) {
if (op == op_kind_t::plus && offset_size == 4 && accum_size == 4)
return { .threads_per_block = 256, .items_per_thread = 14 }; // tuning variant proposes: 512 and 19
...
}
Since we have found that 512 threads per block and 19 items per thread are better, we can update the value in place.
A different case is when we tune beyond what's currently supported by CUB's existing tunings.
This may be because we tune for a new GPU architecture,
in which case a new branch based on the :code:`cuda::compute_capability` passed to the :code:`policy_selector::operator()`
should be introduced, handling this new GPU's compute capability.
Or we tune for new key, value or offset types, etc.,
in which case the existing tuning functions may need additional branches.
There is no general rule on how this extension is done, though.
The implementation may be different for each CUB algorithm.
In the seldom case, that no variant outperforms the baseline,
it must be ensured that any newly added logic correctly falls back to the old tuning values.
There is again no general rule on how this is implemented.
Verification
--------------------------------------------------------------------------------
Once we have selected tunings and implemented them in CUB, we need to verify them.
This process consists of two steps.
Firstly, we need to ensure that adding new tunings and policies did not break existing tunings.
This is most relevant when tunings for new compute capabilities have been added.
To verify this, compile the corresponding benchmarks for the previous compute capabilities
(excluding the new tunings) before and after modifying any tunings,
and compare the generated SASS :code:(`cuobjdump -sass`).
It should not have changed.
Secondly, we must benchmark and compare the performance of the tuned algorithm before and after the tunings have been applied.
This extra step is needed, because the score shown during the tuning analysis is just an aggregated result.
Individual benchmarks may still have regressed for some compile-time workloads.
Fortunately, this is no different than :ref:`running <cub-benchmarking-running>` the corresponding CUB benchmark with and without the changes,
and :ref:`comparing <cub-benchmarking-comparing>` the resulting JSON files.
Such a diff should be supplied to any request to change CUB tunings.
If verification fails for some compile-time workloads (there are regressions), there are two options:
1. Discard the tuning entirely and ensure the tuning selection falls back to the baseline tuning.
2. Narrow down the tuning template specialization to only apply to the workloads where it improves performance,
and fallback where it regressed.
The latter is more complex and may not be justified, if the improvements are small or the use case too narrow.
Use your judgement. Good luck!

22
cccl_upstream/docs/cub/warp_wide.rst Normal file
View File

@@ -0,0 +1,22 @@
.. _warp-module:
Warp-Wide "Collective" Primitives
==================================================
.. toctree::
:glob:
:hidden:
:maxdepth: 2
api/warp
CUB warp-level algorithms are specialized for execution by threads in the same CUDA warp.
These algorithms may only be invoked by ``1 <= n <= 32`` *consecutive* threads in the same warp:
* :cpp:struct:`cub::WarpExchange` rearranges data partitioned across a CUDA warp
* :cpp:class:`cub::WarpLoad` loads a linear segment of items from memory into a CUDA warp
* :cpp:class:`cub::WarpMergeSort` sorts items partitioned across a CUDA warp
* :cpp:struct:`cub::WarpReduce` computes reduction of items partitioned across a CUDA warp
* :cpp:struct:`cub::WarpReduceBatched` computes reduction of multiple batches of items partitioned across a CUDA warp
* :cpp:struct:`cub::WarpScan` computes a prefix scan of items partitioned across a CUDA warp
* :cpp:class:`cub::WarpStore` stores items partitioned across a CUDA warp to a linear segment of memory