Intrinsics#

Bitwise#

__device__ inline int rocprim::get_bit(int x, int i)#

Returns a single bit at ‘i’ from ‘x’.

__device__ inline unsigned int rocprim::bit_count(unsigned int x)#

Bit count.

Returns the number of bit of x set.

__device__ inline unsigned int rocprim::bit_count(unsigned long long x)#

Bit count.

Returns the number of bit of x set.

Warp size#

__host__ __device__ inline constexpr unsigned int rocprim::warp_size()#

[DEPRECATED] Returns a number of threads in a hardware warp.

It is constant for a device. This function is not supported for the gfx1030 architecture and will be removed in a future release. Please use the new host_warp_size() and device_warp_size() functions.

__host__ inline unsigned int rocprim::host_warp_size()#

Returns a number of threads in a hardware warp for the actual device. At host side this constant is available at runtime time only.

It is constant for a device.

__device__ inline constexpr unsigned int rocprim::device_warp_size()#

Returns a number of threads in a hardware warp for the actual target. At device side this constant is available at compile time.

It is constant for a device.

Lane and Warp ID#

Warning

doxygengroup: Cannot find group “intrinsicsmodule_warp_id” in doxygen xml output for project “rocPRIM Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprim/checkouts/docs-5.1.3/docs/.doxygen/docBin/xml

Flat ID#

Warning

doxygengroup: Cannot find group “intrinsicsmodule_flat_id” in doxygen xml output for project “rocPRIM Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprim/checkouts/docs-5.1.3/docs/.doxygen/docBin/xml

Flat Size#

__device__ inline unsigned int rocprim::flat_block_size()#

Returns flat size of a multidimensional block (tile).

__device__ inline unsigned int rocprim::flat_tile_size()#

Returns flat size of a multidimensional tile (block).

Synchronization#

__device__ inline void rocprim::syncthreads()#

Synchronize all threads in a block (tile)

__device__ inline void rocprim::wave_barrier()#

All lanes in a wave come to convergence point simultaneously with SIMT, thus no special instruction is needed in the ISA.

Active threads#

__device__ inline lane_mask_type rocprim::ballot(int predicate)#

Evaluate predicate for all active work-items in the warp and return an integer whose i-th bit is set if and only if predicate is true for the i-th thread of the warp and the i-th thread is active.

Parameters:

predicate – - input to be evaluated for all active lanes

__device__ inline unsigned int rocprim::masked_bit_count(lane_mask_type x, unsigned int add = 0)#

Masked bit count.

For each thread, this function returns the number of active threads which have i-th bit of x set and come before the current thread.