device_wide.rst

Device-Wide Primitives

.. toctree::
   :glob:
   :hidden:
   :maxdepth: 2

   ../api/device

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 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?

• Required: Data types (via template parameters and iterator types) and problem size (num_items)
• Can be nullptr/uninitialized: All input/output pointers (d_in, d_out, etc.)
• Note: The algorithm does not access input data during the query phase

Example pattern:

// Determine temporary storage requirements
void* d_temp_storage = nullptr;
size_t temp_storage_bytes = 0;

cub::DeviceReduce::Sum(
  d_temp_storage, temp_storage_bytes,
  nullptr, nullptr, num_items);  // Input/output pointers can be null

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run the actual algorithm with real pointers
cub::DeviceReduce::Sum(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, num_items);

Single-Phase API (Environment-Based)

Some algorithms provide environment-based overloads that eliminate the two-phase call pattern. These APIs accept an execution environment parameter. See the individual algorithm documentation for availability.

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 withing 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

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::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