alpaka::onAcc::warp Namespace Reference

Functions
template<alpaka::onAcc::concepts::Acc T_Acc>
constexpr auto	activemask (T_Acc const &acc) -> std::conditional_t< T_Acc::getWarpSize()<=32u, uint32_t, uint64_t >
	Return the bit-mask of active lanes for the warp associated with the accelerator.
constexpr bool	all (alpaka::onAcc::concepts::Acc auto const &acc, int32_t predicate)
	Evaluates predicate for all active threads of the warp.
constexpr bool	any (alpaka::onAcc::concepts::Acc auto const &acc, int32_t predicate)
	Evaluates predicate for all active threads of the warp.
template<alpaka::onAcc::concepts::Acc T_Acc>
constexpr auto	ballot (T_Acc const &acc, int32_t predicate) -> std::conditional_t< T_Acc::getWarpSize()<=32u, uint32_t, uint64_t >
	Evaluates predicate for all non-exited threads in a warp and returns a 32- or 64-bit unsigned integer (depending on the accelerator) whose Nth bit is set if and only if predicate evaluates to non-zero for the Nth thread of the warp and the Nth thread is active.
constexpr uint32_t	getLaneIdx (alpaka::onAcc::concepts::Acc auto const &acc)
	Return the lane index of the current thread within its warp.
constexpr uint32_t	getWarpIdx (alpaka::onAcc::concepts::Acc auto const &acc)
	Return the warp index within the block.
template<typename T, alpaka::onAcc::concepts::Acc T_Acc>
constexpr T	shfl (T_Acc const &acc, T const &value, uint32_t srcLane, uint32_t width=getSize< T_Acc >())
	Exchange data between threads within a warp.
template<typename T, alpaka::onAcc::concepts::Acc T_Acc>
constexpr T	shflDown (T_Acc const &acc, T const &value, uint32_t delta, uint32_t width=getSize< T_Acc >())
	Read data from threads with higher lane index within a warp.
template<typename T, alpaka::onAcc::concepts::Acc T_Acc>
constexpr T	shflUp (T_Acc const &acc, T const &value, uint32_t delta, uint32_t width=getSize< T_Acc >())
	Read data from threads with lower lane index within a warp.
template<typename T, alpaka::onAcc::concepts::Acc T_Acc>
constexpr T	shflXor (T_Acc const &acc, T const &value, uint32_t laneMask, uint32_t width=getSize< T_Acc >())
	Exchange data between threads within a warp.
template<concepts::Acc T_Acc>
constexpr uint32_t	getSize ()
	Return the warp size.
template<concepts::Acc T_Acc>
constexpr uint32_t	getSize (T_Acc const &)
	Return the warp size.

Function Documentation

activemask()

template<alpaka::onAcc::concepts::Acc T_Acc>

auto alpaka::onAcc::warp::activemask ( T_Acc const & acc ) -> std::conditional_t< T_Acc::getWarpSize()<=32u, uint32_t, uint64_t >

constexpr

Return the bit-mask of active lanes for the warp associated with the accelerator.

Returns

bit mask where the Nth bit is set to 1 if the corresponding thread is participating the call. The return type can be 64bit or 32bit depending on the API.

Definition at line 33 of file warp.hpp.

all()

bool alpaka::onAcc::warp::all ( alpaka::onAcc::concepts::Acc auto const & acc, int32_t predicate )

constexpr

Evaluates predicate for all active threads of the warp.

It follows the logic of __all_sync(__activemask(), predicate) in CUDA but returns a boolean.

Note:

• The programmer must ensure that all threads calling this function are executing the same line of code. In particular, it is not portable to write if(a) {all} else {all}.

Parameters

predicate

The predicate value for current thread.

Returns

true if all threads predicate non zero, else false

Definition at line 64 of file warp.hpp.

any()

bool alpaka::onAcc::warp::any ( alpaka::onAcc::concepts::Acc auto const & acc, int32_t predicate )

constexpr

Evaluates predicate for all active threads of the warp.

It follows the logic of __any_sync(__activemask(), predicate) in CUDA but returns a boolean.

Note:

• The programmer must ensure that all threads calling this function are executing the same line of code. In particular, it is not portable to write if(a) {any} else {any}.

Parameters

predicate

The predicate value for current thread.

Returns

true if at least one threads predicate is non zero, else false

Definition at line 83 of file warp.hpp.

ballot()

template<alpaka::onAcc::concepts::Acc T_Acc>

auto alpaka::onAcc::warp::ballot ( T_Acc const & acc, int32_t predicate ) -> std::conditional_t< T_Acc::getWarpSize()<=32u, uint32_t, uint64_t >

constexpr

Evaluates predicate for all non-exited threads in a warp and returns a 32- or 64-bit unsigned integer (depending on the accelerator) whose Nth bit is set if and only if predicate evaluates to non-zero for the Nth thread of the warp and the Nth thread is active.

It follows the logic of __ballot_sync(__activemask(), predicate) in CUDA.

Note:

• The programmer must ensure that all threads calling this function are executing the same line of code. In particular, it is not portable to write if(a) {ballot} else {ballot}.

Parameters

predicate

The predicate value for current thread.

Returns

bit mask where the Nth bit is set to 1 if the corresponding threads predicate was non zero. The return type can be 64bit or 32bit depending on the API.

Definition at line 107 of file warp.hpp.

getLaneIdx()

uint32_t alpaka::onAcc::warp::getLaneIdx ( alpaka::onAcc::concepts::Acc auto const & acc )

constexpr

Return the lane index of the current thread within its warp.

Definition at line 41 of file warp.hpp.

getSize() [1/2]

template<concepts::Acc T_Acc>

uint32_t alpaka::onAcc::warp::getSize ( )

constexpr

Return the warp size.

A warp is a collection of threads which work in lock step (executing the same command). The warp size can be larger than the number of threads executed in the kernel/ thread block.

Definition at line 122 of file warp.hpp.

getSize() [2/2]

template<concepts::Acc T_Acc>

uint32_t alpaka::onAcc::warp::getSize ( T_Acc const & )

constexpr

Return the warp size.

A warp is a collection of threads which work in lock step (executing the same command). The warp size can be larger than the number of threads executed in the kernel/ thread block.

Definition at line 128 of file warp.hpp.

getWarpIdx()

uint32_t alpaka::onAcc::warp::getWarpIdx ( alpaka::onAcc::concepts::Acc auto const & acc )

constexpr

Return the warp index within the block.

Definition at line 47 of file warp.hpp.

shfl()

template<typename T, alpaka::onAcc::concepts::Acc T_Acc>

T alpaka::onAcc::warp::shfl ( T_Acc const & acc, T const & value, uint32_t srcLane, uint32_t width = getSize<T_Acc>() )

constexpr

Exchange data between threads within a warp.

Effectively executes:

__shared__ int32_t values[warpsize];
values[threadIdx.x] = value;
__syncthreads();
return values[width*(threadIdx.x/width) + srcLane%width];

However, it does not use shared memory.

Commonly used with width = warpsize (the default), (returns values[srcLane])

This method supports to be called in diverging control flow branches if you only query values from threads within the same branch path.

Parameters

value	value to broadcast, only used if other thread is addressing the lane of this thread.
srcLane	source lane index within the group range [0; width).
width	number of threads receiving a single value, must be a power of 2.

Returns

val from the thread index srcLane.

Definition at line 158 of file warp.hpp.

shflDown()

template<typename T, alpaka::onAcc::concepts::Acc T_Acc>

T alpaka::onAcc::warp::shflDown ( T_Acc const & acc, T const & value, uint32_t delta, uint32_t width = getSize<T_Acc>() )

constexpr

Read data from threads with higher lane index within a warp.

It copies from a lane with higher ID relative to caller. The lane ID is calculated by adding delta to the caller's lane ID.

Effectively executes:

__shared__ int32_t values[warpsize];
values[threadIdx.x] = value;
__syncthreads();
return (threadIdx.x % width + delta < width) ? values[threadIdx.x + delta] : values[threadIdx.x];

However, it does not use shared memory.

Notes:

• The programmer must ensure that all threads calling this function (and the srcLane) are executing the same line of code. In particular it is not portable to write if(a) {shfl} else {shfl}.
• Commonly used with width = warpsize (the default), (returns values[threadIdx.x+delta] if threadIdx.x+delta < warpsize)

Parameters

value	value to broadcast
delta	corresponds to the delta used to compute the lane ID
width	size of the group participating in the shuffle operation, must be a power of 2.

Returns

the value from the thread index lane ID + delta within the group build by width, else value.

Definition at line 194 of file warp.hpp.

shflUp()

template<typename T, alpaka::onAcc::concepts::Acc T_Acc>

T alpaka::onAcc::warp::shflUp ( T_Acc const & acc, T const & value, uint32_t delta, uint32_t width = getSize<T_Acc>() )

constexpr

Read data from threads with lower lane index within a warp.

It copies from a lane with lower ID relative to caller. The lane ID is calculated by subtracting delta from the caller's lane ID.

Effectively executes:

__shared__ int32_t values[warpsize];
values[threadIdx.x] = value;
__syncthreads();
return (threadIdx.x % width >= delta) ? values[threadIdx.x - delta] : values[threadIdx.x];

However, it does not use shared memory.

Notes:

• The programmer must ensure that all threads calling this function (and the srcLane) are executing the same line of code. In particular it is not portable to write if(a) {shfl} else {shfl}.

Commonly used with width = warpsize (the default), (returns values[threadIdx.x - delta] if threadIdx.x >= delta)

Parameters

value	value to broadcast
delta	corresponds to the delta used to compute the lane ID
width	size of the group participating in the shuffle operation, must be a power of 2.

Returns

the value from the thread index lane ID + delta within the group build by width, else value.

Definition at line 229 of file warp.hpp.

shflXor()

template<typename T, alpaka::onAcc::concepts::Acc T_Acc>

T alpaka::onAcc::warp::shflXor ( T_Acc const & acc, T const & value, uint32_t laneMask, uint32_t width = getSize<T_Acc>() )

constexpr

Exchange data between threads within a warp.

It copies from a lane based on bitwise XOR of own lane ID. The lane ID is calculated by performing a bitwise XOR of the caller's lane ID with laneMask

Effectively executes:

__shared__ int32_t values[warpsize];
values[threadIdx.x] = value;
__syncthreads();
int lane = threadIdx.x ^ laneMask;
return values[lane / width > threadIdx.x / width ? threadIdx.x : lane];

However, it does not use shared memory.

Notes:

• The programmer must ensure that all threads calling this function (and the srcLane) are executing the same line of code. In particular it is not portable to write if(a) {shfl} else {shfl}.
• Commonly used with width = warpsize (the default), (returns values[threadIdx.x^laneMask])
• Width must be a power of 2.

  Parameters

  value	value to broadcast
  laneMask	mask applied to the thread lane index within the subgroup created by width.
  width	size of the group participating in the shuffle operation, must be a power of 2.

  Returns

  the value from the thread index lane ID

Definition at line 266 of file warp.hpp.