Exchange#

template<class T, unsigned int BlockSizeX, unsigned int ItemsPerThread, unsigned int BlockSizeY = 1, unsigned int BlockSizeZ = 1>
class block_exchange#

The block_exchange class is a block level parallel primitive which provides methods for rearranging items partitioned across threads in a block.

Overview

  • The block_exchange class supports the following rearrangement methods:

    • Transposing a blocked arrangement to a striped arrangement.

    • Transposing a striped arrangement to a blocked arrangement.

    • Transposing a blocked arrangement to a warp-striped arrangement.

    • Transposing a warp-striped arrangement to a blocked arrangement.

    • Scattering items to a blocked arrangement.

    • Scattering items to a striped arrangement.

  • Data is automatically be padded to ensure zero bank conflicts.

Examples

In the examples exchange operation is performed on block of 128 threads, using type int with 8 items per thread.

__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.blocked_to_striped(items, items, storage);
    ...
}

Template Parameters:
  • T – - the input type.

  • BlockSize – - the number of threads in a block.

  • ItemsPerThread – - the number of items contributed by each thread.

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 type with other storage types to increase shared memory reusability.

Public Functions

template<class U>
__device__ inline void blocked_to_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread])#

Transposes a blocked arrangement of items to a striped arrangement across the thread block.

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

template<class U>
__device__ inline void blocked_to_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], storage_type &storage)#

Transposes a blocked arrangement of items to a striped arrangement across the thread block, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.blocked_to_striped(items, items, storage);
    ...
}

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U>
__device__ inline void striped_to_blocked(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread])#

Transposes a striped arrangement of items to a blocked arrangement across the thread block.

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

template<class U>
__device__ inline void striped_to_blocked(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], storage_type &storage)#

Transposes a striped arrangement of items to a blocked arrangement across the thread block, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.striped_to_blocked(items, items, storage);
    ...
}

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U>
__device__ inline void blocked_to_warp_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread])#

Transposes a blocked arrangement of items to a warp-striped arrangement across the thread block.

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

template<class U>
__device__ inline void blocked_to_warp_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], storage_type &storage)#

Transposes a blocked arrangement of items to a warp-striped arrangement across the thread block, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.blocked_to_warp_striped(items, items, storage);
    ...
}

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U>
__device__ inline void warp_striped_to_blocked(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread])#

Transposes a warp-striped arrangement of items to a blocked arrangement across the thread block.

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

template<class U>
__device__ inline void warp_striped_to_blocked(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], storage_type &storage)#

Transposes a warp-striped arrangement of items to a blocked arrangement across the thread block, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.warp_striped_to_blocked(items, items, storage);
    ...
}

Template Parameters:

U – - [inferred] the output type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U, class Offset>
__device__ inline void scatter_to_blocked(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread])#

Scatters items to a blocked arrangement based on their ranks across the thread block.

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[out] - array that has rank of data.

template<class U, class Offset>
__device__ inline void gather_from_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread])#

Gathers items from a striped arrangement based on their ranks across the thread block.

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[out] - array that has rank of data.

template<class U, class Offset>
__device__ inline void scatter_to_blocked(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread], storage_type &storage)#

Scatters items to a blocked arrangement based on their ranks across the thread block, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    int ranks[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.scatter_to_blocked(items, items, ranks, storage);
    ...
}

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[out] - array that has rank of data.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U, class Offset>
__device__ inline void gather_from_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread], storage_type &storage)#

Gathers items from a striped arrangement based on their ranks across the thread block, using temporary storage.

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[out] - array that has rank of data.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U, class Offset>
__device__ inline void scatter_to_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread])#

Scatters items to a striped arrangement based on their ranks across the thread block.

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[out] - array that has rank of data.

template<class U, class Offset>
__device__ inline void scatter_to_striped(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread], storage_type &storage)#

Scatters items to a striped arrangement based on their ranks across the thread block, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    int ranks[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.scatter_to_striped(items, items, ranks, storage);
    ...
}

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[out] - array that has rank of data.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U, class Offset>
__device__ inline void scatter_to_striped_guarded(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread])#

Scatters items to a striped arrangement based on their ranks across the thread block, guarded by rank.

Overview

  • Items with rank -1 are not scattered.

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[in] - array that has rank of data.

template<class U, class Offset>
__device__ inline void scatter_to_striped_guarded(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread], storage_type &storage)#

Scatters items to a striped arrangement based on their ranks across the thread block, guarded by rank, using temporary storage.

Overview

  • Items with rank -1 are not scattered.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    int ranks[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.scatter_to_striped_guarded(items, items, ranks, storage);
    ...
}

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[in] - array that has rank of data.

  • storage[in] - reference to a temporary storage object of type storage_type.

template<class U, class Offset, class ValidFlag>
__device__ inline void scatter_to_striped_flagged(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread], const ValidFlag (&is_valid)[ItemsPerThread])#

Scatters items to a striped arrangement based on their ranks across the thread block, with a flag to denote validity.

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

  • ValidFlag – - [inferred] the validity flag type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[in] - array that has rank of data.

  • is_valid[in] - array that has flags to denote validity.

template<class U, class Offset, class ValidFlag>
__device__ inline void scatter_to_striped_flagged(const T (&input)[ItemsPerThread], U (&output)[ItemsPerThread], const Offset (&ranks)[ItemsPerThread], const ValidFlag (&is_valid)[ItemsPerThread], storage_type &storage)#

Scatters items to a striped arrangement based on their ranks across the thread block, with a flag to denote validity, using temporary storage.

Storage reusage

Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.
__global__ void example_kernel(...)
{
    // specialize block_exchange for int, block of 128 threads and 8 items per thread
    using block_exchange_int = rocprim::block_exchange<int, 128, 8>;
    // allocate storage in shared memory
    __shared__ block_exchange_int::storage_type storage;

    int items[8];
    int ranks[8];
    int flags[8];
    ...
    block_exchange_int b_exchange;
    b_exchange.scatter_to_striped_flagged(items, items, ranks, flags, storage);
    ...
}

Template Parameters:
  • U – - [inferred] the output type.

  • Offset – - [inferred] the rank type.

  • ValidFlag – - [inferred] the validity flag type.

Parameters:
  • input[in] - array that data is loaded from.

  • output[out] - array that data is loaded to.

  • ranks[in] - array that has rank of data.

  • is_valid[in] - array that has flags to denote validity.

  • storage[in] - reference to a temporary storage object of type storage_type.