[DOC] Add temp_storage_bytes usage guide

cub/examples/device/example_device_reduce.cu

@@ -148,6 +148,7 @@ int main(int argc, char** argv)
   int* d_out = nullptr;
   CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * 1));
 
+  // example-begin temp-storage-query
   // Request and allocate temporary storage
   void* d_temp_storage      = nullptr;
   size_t temp_storage_bytes = 0;
@@ -156,6 +157,7 @@ int main(int argc, char** argv)
 
   // Run
   CubDebugExit(DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items));
+  // example-end temp-storage-query
 
   // Check for correctness (and display results, if specified)
   int compare = CompareDeviceResults(&h_reference, d_out, 1, g_verbose, g_verbose);

docs/cub/Doxyfile

@@ -71,7 +71,7 @@ ALIASES += "blocksize=The number of threads in the block is a multiple of the ar
 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=When ``d_temp_storage`` is ``nullptr``, no work is done and the required allocation size is returned in ``temp_storage_bytes``."
+ALIASES += "devicestorage=When ``d_temp_storage`` is ``nullptr``, no work is done and the required allocation size is returned in ``temp_storage_bytes``. 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."

docs/cub/api_docs/device_wide.rst

@@ -11,6 +11,72 @@ Device-Wide Primitives
    ../api/device
 
 
+.. _device-temp-storage:
+
+Determining Temporary Storage Requirements
+++++++++++++++++++++++++++++++++++++++++++++++++
+
+**Two-Phase API** (Traditional)
+
+Most CUB device-wide algorithms follow a two-phase usage pattern:
+
+1. **Query Phase**: Call the algorithm with ``d_temp_storage = nullptr`` to determine the required temporary storage size
+2. **Execution Phase**: Allocate storage and call the algorithm again to perform the actual operation
+
+**What arguments are needed during the query phase?**
+
+* **Template instantiation**: The query call must use the same template arguments as the execution call.
+* **Argument access**: Aside from ``d_temp_storage``, ``temp_storage_bytes``, and the problem-size arguments, no parameters are accessed during the query phase, so their values may be indeterminate. The dispatch layer returns before launching kernels or touching user storage.
+* **Current device**: The computed temporary storage size is valid only when the execution call 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
+
+**Single-Phase API** (Environment-Based)
+
+Environment-based overloads are rolling out across CUB device-wide primitives. They remove the manual query/execute split by obtaining the temporary storage from a memory resource queried from the execution environment argument.
+
+Key properties of the environment argument:
+
+- It is defaulted and appears as the last argument.
+- Streams can be specified with ``cuda::get_stream`` properties.
+- 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.
+
+Example (centralized control via a single environment argument):
+
+.. code-block:: c++
+
+   #include <cub/device/device_reduce.cuh>
+   #include <cuda/std/execution>
+   #include <cuda/stream_ref>
+   #include <cuda/__memory_resource/get_memory_resource.h>
+   #include <cuda/__execution/determinism.h>
+
+   // Build an execution environment with stream, memory resource, and determinism
+   cudaStream_t stream = /* ... */;
+   auto stream_env = cuda::std::execution::prop{cuda::get_stream_t{}, cuda::stream_ref{stream}};
+
+   auto mr = /* CCCL-provided or user-defined device_memory_resource */;
+   auto mr_env = cuda::std::execution::prop{cuda::mr::__get_memory_resource_t{}, mr};
+
+   auto det_env = cuda::execution::require(cuda::execution::determinism::gpu_to_gpu);
+
+   auto env = cuda::std::execution::env{stream_env, mr_env, det_env};
+
+   // Single-phase API (no explicit temp storage, environment last and defaulted)
+   cub::DeviceReduce::Reduce(d_in, d_out, num_items, cuda::std::plus<>{}, init, env);
+
+The remainder of this page focuses on the traditional two-phase pattern; see individual algorithm documentation for the
+availability and specifics of single-phase overloads.
+
+
 CUB device-level single-problem parallel algorithms:
 
 * :cpp:struct:`cub::DeviceAdjacentDifference` computes the difference between adjacent elements residing within device-accessible memory
@@ -21,7 +87,7 @@ CUB device-level single-problem parallel algorithms:
 * :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 withing a sequence 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