Store#
Class#
-
template<class T, unsigned int ItemsPerThread, unsigned int WarpSize = ::rocprim::device_warp_size(), warp_store_method Method = warp_store_method::warp_store_direct>
class warp_store# The
warp_store
class is a warp level parallel primitive which provides methods for storing an arrangement of items into a blocked/striped arrangement on continous memory.- Overview
The
warp_store
class has a number of different methods to store data:
- Example:
In the example a store operation is performed on a warp of 8 threads, using type
int
and 4 items per thread.__global__ void example_kernel(int * output, ...) { constexpr unsigned int threads_per_block = 128; constexpr unsigned int threads_per_warp = 8; constexpr unsigned int items_per_thread = 4; constexpr unsigned int warps_per_block = threads_per_block / threads_per_warp; const unsigned int warp_id = hipThreadIdx_x / threads_per_warp; const int offset = blockIdx.x * threads_per_block * items_per_thread + warp_id * threads_per_warp * items_per_thread; int items[items_per_thread]; rocprim::warp_store<int, items_per_thread, threads_per_warp, load_method> warp_store; warp_store.store(output + offset, items); ... }
- Template Parameters:
T – - the output/output type.
ItemsPerThread – - the number of items to be processed by each thread.
WarpSize – - the number of threads in a warp. It must be a divisor of the kernel block size.
Method – - the method to store data.
Public Types
-
using storage_type = storage_type_#
Struct used to allocate a temporary memory that is required for thread communication during operations provided by related parallel primitive.
Depending on the implemention the operations exposed by parallel primitive may require a temporary storage for thread communication. The storage should be allocated using keywords
. It can be aliased to an externally allocated memory, or be a part of a union with other storage types to increase shared memory reusability.
Public Functions
-
template<class OutputIterator>
__device__ inline void store(OutputIterator output, T (&items)[ItemsPerThread], storage_type&)# Stores an arrangement of items from across the warp into an arrangement on continuous memory.
- Overview
The type
T
must be such that an object of typeOutputIterator
can be dereferenced and then implicitly assigned fromT
.
- Storage reusage
Synchronization barrier should be placed before
storage
is reused or repurposed:__syncthreads()
orrocprim::syncthreads()
.
- Template Parameters:
OutputIterator – - [inferred] an iterator type for output (can be a simple pointer.
- Parameters:
output – [out] - the output iterator to store to.
items – [in] - array that data is read from.
-
template<class OutputIterator>
__device__ inline void store(OutputIterator output, T (&items)[ItemsPerThread], unsigned int valid, storage_type&)# Stores an arrangement of items from across the warp into an arrangement on continuous memory, which is guarded by range
valid
, using temporary storage.- Overview
The type
T
must be such that an object of typeOutputIterator
can be dereferenced and then implicitly assigned fromT
.
- Storage reusage
Synchronization barrier should be placed before
storage
is reused or repurposed:__syncthreads()
orrocprim::syncthreads()
.
- Template Parameters:
OutputIterator – - [inferred] an iterator type for output (can be a simple pointer.
- Parameters:
output – [out] - the output iterator to store to.
items – [in] - array that data is read from.
valid – [in] - maximum range of valid numbers to read.
Algorithms#
-
enum class rocprim::warp_store_method#
warp_store_method
enumerates the methods available to store a blocked/striped arrangement of items into a blocked/striped arrangement in continuous memoryValues:
-
enumerator warp_store_direct#
A blocked arrangement of items is stored into a blocked arrangement on continuous memory.
- Performance Notes:
Performance decreases with increasing number of items per thread (stride between reads), because of reduced memory coalescing.
-
enumerator warp_store_striped#
A striped arrangement of items is stored into a blocked arrangement on continuous memory.
-
enumerator warp_store_vectorize#
A blocked arrangement of items is stored into a blocked arrangement on continuous memory using vectorization as an optimization.
- Performance Notes:
Performance remains high due to increased memory coalescing, provided that vectorization requirements are fulfilled. Otherwise, performance will default to
warp_store_direct
.
- Requirements:
The output offset (
block_output
) must be quad-item aligned.The following conditions will prevent vectorization and switch to default
warp_store_direct:
ItemsPerThread
is odd.The datatype
T
is not a primitive or a HIP vector type (e.g. int2, int4, etc.
-
enumerator warp_store_transpose#
A blocked arrangement of items is locally transposed and stored as a striped arrangement of data on continuous memory.
- Performance Notes:
Performance remains high due to increased memory coalescing, regardless of the number of items per thread.
Performance may be better compared to
warp_store_direct
andwarp_store_vectorize
due to reordering on local memory.
-
enumerator default_method#
Defaults to
warp_store_direct
.
-
enumerator warp_store_direct#