### 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 // 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 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 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; // 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() ); ... ``` -------------------------------- ### 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; 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 //or // 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 // 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, 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; // 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() ); ... ``` -------------------------------- ### 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; // 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() ); ... } ``` -------------------------------- ### 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; 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 // 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 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 // 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 // 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; // 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() ); ... } ``` -------------------------------- ### 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; 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(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 // 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 keys(keys_input, keys_tmp); rocprim::double_buffer 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 // 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 // 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 // Prepare input and tmp (declare pointers, allocate device memory etc.) size_t input_size; float * input; float * tmp; // Create double-buffer rocprim::double_buffer 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 struct custom_type { int i; double d; }; struct custom_type_decomposer { rocprim::tuple operator()(custom_type& key) const { return rocprim::tuple(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 // 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 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 // 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 struct custom_type { int i; double d; }; struct custom_type_decomposer { rocprim::tuple operator()(custom_type& key) const { return rocprim::tuple(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 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 struct custom_type { int i; double d; }; struct custom_type_decomposer { rocprim::tuple operator()(custom_type& key) const { return rocprim::tuple(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 keys(keys_input, keys_tmp); rocprim::double_buffer 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 struct custom_type { int i; double d; }; struct custom_type_decomposer { rocprim::tuple operator()(custom_type& key) const { return rocprim::tuple(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 keys(keys_input, keys_tmp); rocprim::double_buffer 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 // 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 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(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(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 // 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 class my_alg { private: static constexpr unsigned int wave_size = 32; public: __device__ void some_api(T& a, T& b) { some_opaque_impl(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; // 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() ); ... } ``` -------------------------------- ### 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 // 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; // 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 ); ... } ``` -------------------------------- ### Block Shuffle Down Operation Example Source: https://rocm.docs.amd.com/projects/rocPRIM/en/latest/block_ops/ops_classes/shuffle.html Demonstrates the usage of the block shuffle 'down' operation within a CUDA kernel. This example specializes the block shuffle for integers and a logical warp size of 192 threads. ```C++ __global__ void example_kernel(...) { // specialize block__shuffle_int for int and logical warp of 192 threads using block__shuffle_int = rocprim::block_shuffle; int value = ...; // execute block shuffle block__shuffle_int().down( value, // input value // output ); ... } ```