Utility types

Utility types#

Double buffer#

template<class T> class double_buffer#

This class provides an convenient way to do double buffering.

It tracks two versions of the buffer (“current” and “alternate”), which can be accessed with functions of the same name. You can also swap the buffers.

Public Functions

__host__ __device__ inline double_buffer()#: Constructs an empty double buffer object, initializing the buffer pointers to nullptr.

__host__ __device__ inline double_buffer(T *current, T *alternate)#

Contructs a double buffer object using the supplied buffer pointers.

Parameters:

current – Pointer to the buffer to designate as “current” (in use).
alternate – Pointer to the buffer to designate as “alternate” (not in use)

__host__ __device__ inline T *current() const#: Returns a pointer to the current buffer.

__host__ __device__ inline T *alternate() const#: Returns a pointer to the alternate buffer.

__host__ __device__ inline void swap()#: Swaps the current and alternate buffers.

Future value#

template<typename T, typename Iter = T*> class future_value#

Allows passing values that are not yet known at launch time as parameters to device algorithms.

Note

It is the users responsibility to ensure that value is available when the algorithm executes. This can be guaranteed with stream dependencies or explicit external synchronization.

int* intermediate_result = nullptr;
hipMalloc(reinterpret_cast<void**>(&intermediate_result), sizeof(intermediate_result));
hipLaunchKernelGGL(compute_intermediate, blocks, threads, 0, stream, arg1, arg2, intermediate_result);
const auto initial_value = rocprim::future_value<int>{intermediate_result};
rocprim::exclusive_scan(temporary_storage,
                        storage_size,
                        input,
                        output,
                        initial_value,
                        size);
hipFree(intermediate_result)

Template Parameters:

T –
Iter –

Public Types

using value_type = T #: The type of the value this future will store.

using iterator_type = Iter #: An iterator type that can point at the value_type.

Public Functions

__host__ __device__ inline explicit future_value(const Iter iter)#

Constructs a future value.

Parameters:: iter – An iterator that will point to the value when it becomes available.

__host__ __device__ inline operator T()#: Returns the value by dereferencing the iterator that the constructor was passed.

Note

The value must be available at the point this is called.

__host__ __device__ inline operator T() const#: Returns the value by dereferencing the iterator that the constructor was passed.

Note

The value must be available at the point this is called.

Key-value pair#

template<class Key_, class Value_> struct key_value_pair#

Convenience struct that allows you to store key-value pairs as a composite entity.

Public Types

using key_type = Key_#: the key’s type

using value_type = Value_#: the value’s type

Public Functions

__host__ __device__ inline key_value_pair(const key_type key, const value_type value)#: Constructs a key-value pair using the supplied values.

__host__ __device__ inline bool operator!=(const key_value_pair &kvb)#: Returns true if the given key-value pair is not equal to this one. Otherwise returns false.

Public Members

key_type key#: the stored key

value_type value#: the stored value

Tuple#

template<class ...Types> class tuple#

Fixed-size collection of heterogeneous values.

Uninitialized Array#

template<typename T, unsigned int Count, size_t Alignment = alignof(T)> class uninitialized_array#

Provides indexed & typed access to uninitialized memory. To be used with ROCPRIM_SHARED_MEMORY

Note

This class should be used to ensure that writes to the uninitialized memory block occur only via placement new.

Note

This class is non-copyable.

Note

The recommended pattern for usage is to first fill the array via calls to emplace, then read-only via get_unsafe_array(). Writing to the array reference returned by get_unsafe_array() should be avoided, if possible.

Note

The value of Alignment MUST be a valid alignment for T.

Template Parameters:

T – The item type which is provided via the accessors.
Count – The number of T items to store.
Alignment – The alignment of the backing storage, in bytes.

Public Types

using c_array_t = T[Count]#: Type of the represented C array.

Public Functions

__host__ __device__ uninitialized_array(uninitialized_array&&) = default#: Default move constructor.

__host__ __device__ uninitialized_array &operator=(uninitialized_array&&) = default#: Default move assignment.

template<typename ...Args> __host__ __device__ inline T &emplace(const unsigned int index, Args&&... args)#

Constructs a value in-place at the specified array index.

Note

This function calls the constructor of T with the specified arguments. If an instance of T& is passed, the copy constructor is called, whereas in the case of a T&&, the move constructor is called.

Note

If an object is created at the same index more than once, the old object’s destructor is not called. If T is not trivially destructible, the behaviour is undefined.

Template Parameters:

Args – The types of the argument values.

Parameters:

index – The index in the array where the new object is constructed.
args – The arguments to call the constructor with.

Returns:

A reference to the newly constructed element.

__host__ __device__ inline c_array_t &get_unsafe_array()#: Returns a reference to the underlying memory as a typed array.

Note

Manipulating items in the returned array reference at indices that were not previously filled by calls to emplace MUST be avoided.

Utility types

Contents

Utility types#

Double buffer#

Future value#

Key-value pair#

Tuple#

Uninitialized Array#