rocPRIM

https://rocm.docs.amd.com/projects/rocPRIM/en/latest/
Admin
rocPRIM is a header-only library that provides HIP parallel primitives. The purpose of the library...
Tokens:301,587
Snippets:1,945
Trust Score:10
Update:3 months ago
Show doc for...
Tokens:
Raw
### Build Examples, Tests, and Benchmarks

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/install/rocPRIM-build-install-linux.html

Build and install rocPRIM along with its examples, tests, and benchmarks. This involves enabling tests with CMake and then running make and make install.

```bash
export CXX=hipcc
cmake -DBUILD_TEST=ON ../.
make
sudo make install
```

--------------------------------

### Device-level nth_element Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/nth_element.html

Demonstrates performing a device-level nth_element operation on an array of unsigned integers. It shows how to get the required temporary storage size, allocate it, and then execute the nth_element function.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
size_t nth;                 // e.g., 4
unsigned int * keys;  // e.g., [ 6, 3, 5, 4, 1, 8, 2, 7 ]

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, nth, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform nth_element
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, nth, input_size
);
// possible keys:   [ 1, 3, 4, 2, 5, 8, 7, 6 ]

```

--------------------------------

### Device-Level Two-Way Partition Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/partition.html

Demonstrates partitioning an array of integers into even (selected) and odd (rejected) numbers using the rocPrim partition_two_way primitive. It shows how to get the temporary storage size, allocate it, and then perform the partition.

```cpp
#include <rocprim/rocprim.hpp>

auto predicate = 
    [] (int a) -> bool 
    { 
        return (a%2) == 0; 
    };

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;     // e.g., 8
int * input;           // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
int * selected_output; // empty array of at least 4 elements
int * rejected_output; // empty array of at least 4 elements
size_t * output_count; // empty array of 1 element

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::partition_two_way(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input,
    selected_output,
    rejected_output,
    selected_output_count,
    input_size,
    predicate
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform partition
rocprim::partition_two_way(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input,
    selected_output,
    rejected_output,
    selected_output_count,
    input_size,
    predicate
);
// output_selected: [2, 4, 6, 8]
// output_rejected: [1, 3, 5, 7]
// selected_output_count: 4

```

--------------------------------

### Install rocPRIM on Windows

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/install/rocPRIM-build-install-windows.html

Run this command in the cloned rocPRIM directory to install rocPRIM to C:\hipSDK\include\.

```bash
cd rocPRIM

python3 rmake.py -i

```

--------------------------------

### Device-Level Three-Way Partition Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/partition.html

Demonstrates how to perform a device-level three-way partition on an array of integers. Even numbers go to the first partition, odd and 3-divisible numbers go to the second, and the rest are unselected. This snippet shows the process of getting the required temporary storage size, allocating it, and then executing the partition.

```cpp
#include <rocprim/rocprim.hpp>

auto first_predicate = 
    [] (int a) -> bool 
    {
        return (a%2) == 0;
    };
auto second_predicate = 
    [] (int a) -> bool 
    {
        return (a%3) == 0;
    };

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
int * input;                // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
int * output_first_part;    // array of 8 elements
int * output_second_part;   // array of 8 elements
int * output_unselected;    // array of 8 elements
size_t * output_count;      // array of 2 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::partition_three_way(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input,
    output_first_part, output_second_part, output_unselected,
    output_count,
    input_size,
    first_predicate,
    second_predicate
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform partition
rocprim::partition_three_way(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input,
    output_first_part, output_second_part, output_unselected,
    output_count,
    input_size,
    first_predicate,
    second_predicate
);
// elements denoted by '*' were not modified
// output_first_part:  [2, 4, 6, 8, *, *, *, *]
// output_second_part: [3, *, *, *, *, *, *, *]
// output_unselected:  [1, 5, 7, *, *, *, *, *]
// output_count:       [4, 1]

```

--------------------------------

### Build and Install rocPRIM

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/install/rocPRIM-build-install-linux.html

Configure, build, and install the rocPRIM library using CMake and make. This command builds the library without examples, tests, or benchmarks.

```bash
cmake ../.
make install
```

--------------------------------

### Build rocPRIM Examples, Tests, and Benchmarks

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/install/rocPRIM-build-install-windows.html

Use the -c option with rmake.py to build the examples, tests, and benchmarks for rocPRIM.

```bash
python3 rmake.py -c
```

--------------------------------

### Example Kernel for Inclusive Scan

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/scan.html

Demonstrates an example kernel performing an inclusive maximum scan on a block of 128 threads, each with two long values. Shows specialization, storage allocation, and invocation of the scan operation.

```cpp
__global__ void example_kernel(...) // blockDim.x = 128
{
    // specialize block_scan for long and block of 128 threads
    using block_scan_long = rocprim::block_scan<long, 128>;
    // allocate storage in shared memory for the block
    __shared__ block_scan_long::storage_type storage;

    long input[2] = ...;
    long output[2];
    long reduction;
    // execute inclusive min scan
    block_scan_long().inclusive_scan(
        input,
        output,
        reduction,
        storage,
        rocprim::maximum<long>()
    );
    ...
```

--------------------------------

### Warp Sort Example with Temporary Storage

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/warp_ops/sort.html

Example kernel demonstrating the usage of warp_sort with temporary storage for sorting integer values within a warp.

```cpp
__global__ void example_kernel(...)
{
    int value = ...;
    using warp_sort_int = rp::warp_sort<int, 64>;
    warp_sort_int wsort;
    __shared__ typename warp_sort_int::storage_type storage;
    wsort.sort(value, storage);
    ...
```

--------------------------------

### Adjacent Difference Operation Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/adjacent_difference.html

Demonstrates performing a device-level adjacent difference operation on integer values using a custom binary function. It includes steps for getting the required temporary storage size, allocating it, and then executing the operation.

```cpp
#include <rocprim/rocprim.hpp> //or <rocprim/device/device_adjacent_difference.hpp>

// custom binary function
auto binary_op =
    [] (int a, int b) -> int
    {
        return a - b;
    };

// Prepare input and output (declare pointers, allocate device memory etc.)
std::size_t size; // e.g., 8
int* input1; // e.g., [8, 7, 6, 5, 4, 3, 2, 1]
int* output; // empty array of 8 elements

std::size_t temporary_storage_size_bytes;
void* temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::adjacent_difference(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, size, binary_op
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform adjacent difference
rocprim::adjacent_difference(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, size, binary_op
);
// output: [8, 1, 1, 1, 1, 1, 1, 1]

```

--------------------------------

### Device Memory Copy Example using rocPrim batch_memcpy

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/memcpy.html

This example demonstrates how to use rocPrim's batch_memcpy to copy multiple sections of data from one device memory location to another. It includes calculating temporary storage, allocating it, and then performing the batched copy.

```cpp
#include <rocprim/rocprim.hpp>

// Device allocated data:
int* a;             // e.g, [9, 1, 2, 3, 4, 5, 6, 7, 8]
int* b;             // e.g, [0, 0, 0, 0, 0, 0, 0, 0, 0]

// Batch memcpy parameters:
int   num_copies;   // Number of buffers to copy.
                     // e.g, 4.
int** sources;       // Pointer to source pointers.
                     // e.g, [&a[0], &a[4] &a[7]]
int** destinations;  // Pointer to destination pointers.
                     // e.g, [&b[5], &b[2] &b[0]]
int*  sizes;         // Size of buffers to copy.
                     // e.g., [3 * sizeof(int), 2 * sizeof(int), 2 * sizeof(int)]

// Calculate the required temporary storage.
size_t temporary_storage_size_bytes;
void* temporary_storage_ptr = nullptr;
rocprim::batch_memcpy(
    temporary_storage_ptr,
    temporary_storage_size_bytes,
    sources,
    destinations,
    sizes,
    num_buffers);

// Allocate temporary storage.
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// Copy buffers.
rocprim::batch_memcpy(
    temporary_storage_ptr,
    temporary_storage_size_bytes,
    sources,
    destinations,
    sizes,
    num_copies);

// b is now: [7, 8, 4, 5, 0, 9, 1, 2, 0]
//   3rd copy ^--^  ^--^     ^--^--^ 1st copy
//                2nd copy


```

--------------------------------

### Inclusive and Exclusive Scan Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/warp_ops/scan.html

Demonstrates performing both inclusive and exclusive scan operations across threads in a logical warp. Synchronization barriers like `__syncthreads()` or `rocprim::syncthreads()` are required before reusing storage. This example uses a virtual wave size less than or equal to the hardware wavefront size.

```cpp
template<class BinaryFunction = ::rocprim::plus<T>, unsigned int FunctionVirtualWaveSize = VirtualWaveSize> 
__device__ inline auto scan(T input, T &inclusive_output, T &exclusive_output, T init, storage_type &storage, BinaryFunction scan_op = BinaryFunction()) -> typename std::enable_if<(FunctionVirtualWaveSize <= arch::wavefront::max_size()), void>::type
```

```cpp
__global__ void example_kernel(...) // blockDim.x = 256
{
    // specialize warp_scan for float and logical warp of 32 threads
    using warp_scan_f = rocprim::warp_scan<float, 32>;
    // allocate storage in shared memory
    __shared__ warp_scan_float::storage_type temp[8]; // 256/32 = 8

    int logical_warp_id = threadIdx.x/32;
    float input = ...;
    float ex_output, in_output;
    // execute exclusive min scan
    warp_scan_float().scan(
        input,
        in_output,
        ex_output,
        100.0f, // init
        temp[logical_warp_id],
        rocprim::minimum<float>()
    );
    ...
```

--------------------------------

### Block Reduce Example Kernel

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/reduce.html

Example of using `rocprim::block_reduce` within a CUDA kernel for maximum reduction on `long` values. It demonstrates specialization, storage allocation, and calling the reduce function with a custom operation.

```cpp
__global__ void example_kernel(...) // blockDim.x = 128
{
    // specialize block_reduce for long and block of 128 threads
    using block_reduce_f = rocprim::block_reduce<long, 128>;
    // allocate storage in shared memory for the block
    __shared__ block_reduce_long::storage_type storage;

    long input[2] = ...;
    long output[2];
    // execute max reduce
    block_reduce_long().reduce(
        input,
        output,
        storage,
        rocprim::maximum<long>()
    );
    ...
}
```

--------------------------------

### Example Usage of block_store

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/store.html

Demonstrates how to instantiate and use the block_store operation within a CUDA kernel. Ensure that shared memory for storage is declared and initialized appropriately.

```cuda
__global__ void kernel(...)
{
    int items[8];
    using block_store_int = rocprim::block_store<int, 128, 8>;
    block_store_int bstore;
    __shared__ typename block_store_int::storage_type storage;
    bstore.store(..., items, valid, storage);
    ...
}

```

--------------------------------

### Device-Level Ascending Merge Sort Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Demonstrates performing a device-level ascending merge sort on unsigned integer keys and double values. It shows how to obtain the required temporary storage size and then perform the sort.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
unsigned int * keys_input;  // e.g., [ 6, 3,  5, 4,  1,  8,  2, 7]
double * values_input;      // e.g., [-5, 2, -4, 3, -1, -8, -2, 7]
unsigned int * keys_output; // empty array of 8 elements
double * values_output;     // empty array of 8 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::merge_sort(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, values_input, values_output,
    input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::merge_sort(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, values_input, values_output,
    input_size
);
// keys_output:   [ 1,  2, 3, 4,  5,  6, 7,  8]
// values_output: [-1, -2, 2, 3, -4, -5, 7, -8]

```

--------------------------------

### Example Kernel for Block Load

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/load.html

Demonstrates a basic usage of rocprim::block_load within a CUDA kernel. This example initializes a block_load object and performs a load operation from global memory into a local array. Ensure that the load_method template parameter is correctly specified.

```cpp
__global__ void example_kernel(int * input, ...)
{
    const int offset = blockIdx.x * 128 * 8;
    int items[8];
    rocprim::block_load<int, 128, 8, load_method> blockload;
    blockload.load(input + offset, items);
    ...
}
```

--------------------------------

### Device-Level Search with Unsigned Integers

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/search.html

Performs a device-level search for a sequence of unsigned integers. This example demonstrates how to allocate temporary storage and execute the search operation.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t size;           // e.g., 10
size_t key_size;       // e.g., 3
unsigned int * input;  // e.g., [ 6, 3, 5, 4, 1, 8, 2, 5, 4, 1 ]
unsigned int * key;    // e.g., [ 5, 4, 1 ]
unsigned int * output; // e.g., empty array of size 1

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::search(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, key, output, size, key_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform search
rocprim::search(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, key, output, size, key_size
);
// output:   [ 2 ]

```

--------------------------------

### Device Segmented Radix Sort Keys Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

This example demonstrates how to perform a device-level ascending radix sort on an array of float values using rocprim::segmented_radix_sort_keys. It includes steps for preparing input/output, allocating temporary storage, and executing the sort.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;      // e.g., 8
float * input;          // e.g., [0.6, 0.3, 0.65, 0.4, 0.2, 0.08, 1, 0.7]
float * output;         // empty array of 8 elements
unsigned int segments;  // e.g., 3
int * offsets;          // e.g. [0, 2, 3, 8]

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::segmented_radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size,
    segments, offsets, offsets + 1
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::segmented_radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size,
    segments, offsets, offsets + 1
);
// keys_output: [0.3, 0.6, 0.65, 0.08, 0.2, 0.4, 0.7, 1]

```

--------------------------------

### Example: Inclusive Min Scan with Custom Operator

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/warp_ops/scan.html

Illustrates an inclusive scan operation using a custom binary function (rocprim::minimum) for floating-point values. This example specializes warp_scan for 'float' and a logical warp size of 32, utilizing shared memory for storage. Remember to synchronize threads if reusing storage.

```cpp
__global__ void example_kernel(...) // blockDim.x = 256
{
    // specialize warp_scan for float and logical warp of 32 threads
    using warp_scan_f = rocprim::warp_scan<float, 32>;
    // allocate storage in shared memory
    __shared__ warp_scan_float::storage_type temp[8]; // 256/32 = 8

    int logical_warp_id = threadIdx.x/32;
    float value = ...;
    // execute inclusive min scan
    warp_scan_float().inclusive_scan(
        value, // input
        value, // output
        temp[logical_warp_id],
        rocprim::minimum<float>()
    );
    ...
}

```

--------------------------------

### Get Device Architecture Configuration (Host)

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/reference/developer.html

Selects configuration parameters at runtime based on device architecture determined from the HIP stream.

```cpp
using config = wrapped_ALGO_config<config, Type>;

detail::target_arch target_arch;
hipError_t          result = host_target_arch(stream, target_arch);
if(result != hipSuccess)
{
    return result;
}
const ALGO_config_params params = dispatch_target_arch<config>(target_arch);

```

--------------------------------

### Device Ascending Radix Sort of (Key, Value) Pairs

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

This example demonstrates a device-level ascending radix sort using rocprim::radix_sort_pairs. It sorts an array of unsigned integers (keys) and corresponding doubles (values). The example shows how to prepare input and temporary buffers, obtain the required temporary storage size, allocate it, and then perform the sort. It also illustrates how to optimize performance by specifying the relevant bit range for keys.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and tmp (declare pointers, allocate device memory etc.)
size_t input_size;
unsigned int * keys_input;  // e.g., [ 6, 3,  5, 4,  1,  8,  1, 7]
double * values_input;      // e.g., [-5, 2, -4, 3, -1, -8, -2, 7]
unsigned int * keys_tmp;    // empty array of 8 elements
double*  values_tmp;        // empty array of 8 elements
// Create double-buffers
rocprim::double_buffer<unsigned int> keys(keys_input, keys_tmp);
rocprim::double_buffer<double> values(values_input, values_tmp);

// Keys are in range [0; 8], so we can limit compared bit to bits on indexes
// 0, 1, 2, 3, and 4. In order to do this begin_bit is set to 0 and end_bit
// is set to 5.

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, values, input_size,
    0, 5
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, values, input_size,
    0, 5
);
// keys.current():   [ 1,  1, 3, 4,  5,  6, 7,  8]
// values.current(): [-1, -2, 2, 3, -4, -5, 7, -8]

```

--------------------------------

### Ascending Radix Sort Keys Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Demonstrates how to perform a device-level ascending radix sort on an array of float values using rocprim::radix_sort_keys. It includes steps for preparing input/output, allocating temporary storage, and executing the sort.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;      // e.g., 8
float * input;          // e.g., [0.6, 0.3, 0.65, 0.4, 0.2, 0.08, 1, 0.7]
float * output;         // empty array of 8 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size
);
// keys_output: [0.08, 0.2, 0.3, 0.4, 0.6, 0.65, 0.7, 1]

```

--------------------------------

### Device Descending Radix Sort Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Performs a device-level descending radix sort on an array of integer values. Requires device memory allocation for input, output, and temporary storage.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;    // e.g., 8
int * input;          // e.g., [6, 3, 5, 4, 2, 8, 1, 7]
int * output;         // empty array of 8 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_keys_desc(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_keys_desc(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size
);
// keys_output: [8, 7, 6, 5, 4, 3, 2, 1]

```

--------------------------------

### Ascending Radix Sort Keys Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Demonstrates how to perform a device-level ascending radix sort on an array of float values using rocprim::radix_sort_keys. It includes steps for preparing input data, allocating temporary storage, and executing the sort.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and tmp (declare pointers, allocate device memory etc.)
size_t input_size;
float * input;
float * tmp;
// Create double-buffer
rocprim::double_buffer<float> keys(input, tmp);

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, input_size
);
// keys.current(): [0.08, 0.2, 0.3, 0.4, 0.6, 0.65, 0.7, 1]

```

--------------------------------

### Ascending Radix Sort with Custom Type and Decomposer

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Demonstrates a device-level ascending radix sort on an array of custom types using a custom decomposer. It shows how to get the required temporary storage size, allocate it, and then perform the sort.

```cpp
#include <rocprim/rocprim.hpp>

struct custom_type
{
    int i;
    double d;
};

struct custom_type_decomposer
{
    rocprim::tuple<int&, double&> operator()(custom_type& key) const
    {
        return rocprim::tuple<int&, double&>(key.i, key.d);
    }
};

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;      // e.g., 8
custom_type * input;    // e.g., [{2, 0.6}, {-3, 0.3}, {2, 0.65}, {0, 0.4}, {0, 0.2}, {11, 0.08}, {11, 1}, {-1, 0.7}]
custom_type * output;   // empty array of 8 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size, custom_type_decomposer{} 
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size, custom_type_decomposer{} 
);
// keys_output: [{-3, 0.3}, {-1, 0.7}, {0, 0.2}, {0, 0.4}, {2, 0.6}, {2, 0.65}, {11, 0.08}, {11, 1.0}]

```

--------------------------------

### Device Descending Radix Sort Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Demonstrates how to perform a device-level descending radix sort on an array of integer values using rocprim::radix_sort_keys_desc. It includes steps for preparing input and temporary storage, obtaining the required temporary storage size, allocating it, and finally performing the sort.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and tmp (declare pointers, allocate device memory etc.)
size_t input_size;  // e.g., 8
int * input;        // e.g., [6, 3, 5, 4, 2, 8, 1, 7]
int * tmp;          // empty array of 8 elements
// Create double-buffer
rocprim::double_buffer<int> keys(input, tmp);

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_keys_desc(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_keys_desc(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, input_size
);
// keys.current(): [8, 7, 6, 5, 4, 3, 2, 1]

```

--------------------------------

### Ascending Radix Sort Pairs

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Performs a device-level ascending radix sort of (key, value) pairs. The `Key` type must be arithmetic. The function can be used to get the required temporary storage size by passing `nullptr` for `temporary_storage`. The `begin_bit` and `end_bit` parameters can optimize performance for integer keys within a known range.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
unsigned int * keys_input;  // e.g., [ 6, 3,  5, 4,  1,  8,  1, 7]
double * values_input;      // e.g., [-5, 2, -4, 3, -1, -8, -2, 7]
unsigned int * keys_output; // empty array of 8 elements
double * values_output;     // empty array of 8 elements

// Keys are in range [0; 8], so we can limit compared bit to bits on indexes
// 0, 1, 2, 3, and 4. In order to do this begin_bit is set to 0 and end_bit
// is set to 5.

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, values_input, values_output,
    input_size, 0, 5
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, values_input, values_output,
    input_size, 0, 5
);
// keys_output:   [ 1,  1, 3, 4,  5,  6, 7,  8]
// values_output: [-1, -2, 2, 3, -4, -5, 7, -8]

```

--------------------------------

### Device Ascending Radix Sort with Custom Type and Decomposer

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Performs a device-level ascending radix sort on an array of custom types. Requires defining a struct for the custom type and a decomposer functor to extract sortable members. The example demonstrates allocating temporary storage and executing the sort.

```cpp
#include <rocprim/rocprim.hpp>

struct custom_type
{
    int i;
    double d;
};

struct custom_type_decomposer
{
    rocprim::tuple<int&, double&> operator()(custom_type& key) const
    {
        return rocprim::tuple<int&, double&>(key.i, key.d);
    }
};

// Prepare input and tmp (declare pointers, allocate device memory etc.)
constexpr unsigned int begin_bit = 0;
constexpr unsigned int end_bit = 96;
size_t input_size;   // e.g., 8
custom_type * input; // e.g., [{2, 0.6}, {-3, 0.3}, {2, 0.65}, {0, 0.4}, {0, 0.2}, {11, 0.08}, {11, 1}, {-1, 0.7}]
custom_type * tmp;   // empty array of 8 elements
// Create double-buffer
rocprim::double_buffer<custom_type> keys(input, tmp);

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, input_size, custom_type_decomposer{}, begin_bit, end_bit
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_keys(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, input_size, custom_type_decomposer{}, begin_bit, end_bit
);
// keys.current(): [{-3, 0.3}, {-1, 0.7}, {0, 0.2}, {0, 0.4}, {2, 0.6}, {2, 0.65}, {11, 0.08}, {11, 1.0}]

```

--------------------------------

### Device Descending Radix Sort with Custom Key Type

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Performs a device-level descending radix sort on input keys of a custom type and values of `double`. Requires setup of input/temporary buffers and a decomposer functor to extract sortable fields from the custom key.

```cpp
#include <rocprim/rocprim.h>

struct custom_type
{
    int i;
    double d;
};

struct custom_type_decomposer
{
    rocprim::tuple<int&, double&> operator()(custom_type& key) const
    {
        return rocprim::tuple<int&, double&>(key.i, key.d);
    }
};

// Prepare input and tmp (declare pointers, allocate device memory etc.)
size_t input_size;        // e.g., 8
custom_type * keys_input; // e.g., [{2, 0.6}, {0, 0.3}, {2, 0.65}, {0, 0.4}, {0, 0.2}, {11, 0.08}, {11, 1.0}, {5, 0.7}]
double * values_input;    // e.g., [-5, 2, -4, 3, -1, -8, -2, 7]
custom_type * keys_tmp;   // empty array of 8 elements
double * values_tmp;      // empty array of 8 elements
// Create double-buffers
rocprim::double_buffer<custom_type> keys(keys_input, keys_tmp);
rocprim::double_buffer<double> values(values_input, values_tmp);

// The integer field of the keys is in range 0-11, which can be represented on 4 bits, 
// while for the double member we must specify full bit range [0; 63]. Therefore begin_bit
// is set to 0 and end_bit is set to 68.
constexpr unsigned int begin_bit = 0;
constexpr unsigned int end_bit = 68;

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_pairs_desc(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, values, input_size, custom_type_decomposer{}, begin_bit, end_bit
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_pairs_desc(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, values, input_size, custom_type_decomposer{}, begin_bit, end_bit
);
// keys.current():   [{11, 1.0}, {11, 0.08}, {5, 0.7}, {2, 0.65}, {2, 0.6}, {0, 0.4}, {0, 0.3}, {0, 0.2}]
// values.current(): [-2, -1, 2, 3, -4, -5, 7, -8]

```

--------------------------------

### Device Ascending Radix Sort with Custom Key Type

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/sort.html

Performs a device-level ascending radix sort on custom key types and associated double values. Requires setup of input, temporary storage, and a custom decomposer. The sort operates on specified bit ranges of the key's components.

```cpp
#include <rocprim/rocprim.hpp>

struct custom_type
{
    int i;
    double d;
};

struct custom_type_decomposer
{
    rocprim::tuple<int&, double&> operator()(custom_type& key) const
    {
        return rocprim::tuple<int&, double&>(key.i, key.d);
    }
};

// Prepare input and tmp (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
custom_type * keys_input;   // e.g., [{2, 0.6}, {0, 0.3}, {2, 0.65}, {0, 0.4}, {0, 0.2}, {11, 0.08}, {11, 1.0}, {5, 0.7}]
double * values_input;      // e.g., [-5, 2, -4, 3, -1, -8, -2, 7]
custom_type * keys_tmp;     // empty array of 8 elements
double*  values_tmp;        // empty array of 8 elements
// Create double-buffers
rocprim::double_buffer<custom_type> keys(keys_input, keys_tmp);
rocprim::double_buffer<double> values(values_input, values_tmp);

// The integer field of the keys is in range 0-11, which can be represented on 4 bits, 
// while for the double member we must specify full bit range [0; 63]. Therefore begin_bit
// is set to 0 and end_bit is set to 68.
constexpr unsigned int begin_bit = 0;
constexpr unsigned int end_bit = 68;

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::radix_sort_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, values, input_size, custom_type_decomposer{}, begin_bit, end_bit
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform sort
rocprim::radix_sort_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, values, input_size, custom_type_decomposer{}, begin_bit, end_bit
);
// keys.current():   [{0, 0.2}, {0, 0.3}, {0, 0.4}, {2, 0.6}, {2, 0.65}, {5, 0.7}, {11, 0.08}, {11, 1.0}]
// values.current(): [-1, 2, 3, -5, -4, 7, -8, -2]

```

--------------------------------

### Device Binary Search Example

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/binary_search.html

Demonstrates performing a device-level binary search on an array of integers. It includes steps for preparing input/output, calculating temporary storage, allocating it, and executing the search.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.).
size_t          haystack_size;    // e.g. 10
int *           haystack;         // e.g. {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}
size_t          needles_size;     // e.g. 8
int *           needles;          // e.g. {0, 2, 12, 4, 14, 6, 8, 10}
compare_op_type compare_op;       // e.g. compare_op_type = rocprim::less<int>
size_t *        output;           // empty array of needles_size elements

// Get required size of the temporary storage.
void * temporary_storage = nullptr;
size_t temporary_storage_bytes;
rocprim::binary_search<config>(temporary_storage,
                               temporary_storage_bytes,
                               haystack,
                               needles,
                               output,
                               haystack_size,
                               needles_size,
                               compare_op,
                               stream,
                               debug_synchronous);

// Allocate temporary storage.
hipMalloc(&temporary_storage, temporary_storage_bytes);

// Perform binary search.
rocprim::binary_search<config>(temporary_storage,
                               temporary_storage_bytes,
                               haystack,
                               needles,
                               output,
                               haystack_size,
                               needles_size,
                               compare_op,
                               stream,
                               debug_synchronous);

// output = {1, 1, 0, 1, 0, 1, 1, 0}

```

--------------------------------

### Create Build Directory for Visual Studio Projects

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/install/rocPRIM-build-install-windows.html

Navigate to the 'benchmark' directory and create a 'build' directory to prepare for generating Visual Studio projects.

```bash
cd benchmark
mkdir build

```

--------------------------------

### Device-Level Partial Sort Copy

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/partial_sort.html

Performs a device-level partial sort copy on an array of unsigned integers. This example demonstrates how to get the required temporary storage size, allocate it, and then execute the partial sort.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;
size_t middle;
unsigned int * keys_input;
unsigned int * keys_output;

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::partial_sort_copy(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, middle, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform partial_sort
rocprim::partial_sort_copy(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, middle, input_size
);
// possible keys_output:   [ 1, 2, 3, 4, 5, 9, 9, 9 ]

```

--------------------------------

### Static Wavefront Size Algorithm

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/how-to/rocPRIM-spir-v.html

This example shows a basic algorithm with a statically defined wavefront size. This approach is suitable when the target hardware's wavefront size is known and fixed.

```cpp
template<typename T>
class my_alg
{
private:
    static constexpr unsigned int wave_size = 32;
public:
    __device__ void some_api(T& a, T& b)
    {
        some_opaque_impl<wave_size>(a, b);
    }
};

```

--------------------------------

### Example: Min Reduce Operation with Block Reduce

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/reduce.html

Demonstrates a min reduce operation on a block of 256 threads, each providing a float value. Requires explicit storage allocation and synchronization.

```c++
__global__ void example_kernel(...) // blockDim.x = 256
{
    // specialize block_reduce for float and block of 256 threads
    using block_reduce_f = rocprim::block_reduce<float, 256>;
    // allocate storage in shared memory for the block
    __shared__ block_reduce_float::storage_type storage;

    float input = ...;
    unsigned int valid_items = 250;
    float output;
    // execute min reduce
    block_reduce_float().reduce(
        input,
        output,
        valid_items,
        storage,
        rocprim::minimum<float>()
    );
    ...
}

```

--------------------------------

### Create and Navigate to Build Directory

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/install/rocPRIM-build-install-linux.html

Create a build directory and change into it before configuring the CMake build.

```bash
mkdir build
cd build
```

--------------------------------

### Get Required Temporary Storage Size for nth_element

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/device_ops/nth_element.html

Call `rocprim::nth_element` with a null pointer for temporary storage to get the required allocation size in bytes. This is useful for pre-allocating memory.

```cpp
#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
size_t nth;                 // e.g., 4
unsigned int * keys_input;  // e.g., [ 6, 3, 5, 4, 1, 8, 2, 7 ]
unsigned int * keys_output; // empty array of 8 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, nth, input_size
);
```

--------------------------------

### Example: Block Histogram Kernel

Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/histogram.html

An example CUDA kernel demonstrating the usage of `rocprim::block_histogram` for an integer histogram. It shows specialization, shared memory allocation, and the histogram operation within a kernel.

```cpp
__global__ void example_kernel(...)
{
    // specialize block_histogram for int, logical block of 192 threads,
    // 2 items per thread and a bin size of 192.
    using block_histogram_int = rocprim::block_histogram<int, 192, 2, 192>;
    // allocate storage in shared memory
    __shared__ block_histogram_int::storage_type storage;
    __shared__ int hist[192];

    int value[2];
    ...
    // execute histogram
    block_histogram_int().histogram(
        value, // input
        hist, // output
        storage
    );
    ...
}
```
Tokens:
Raw