asc_ballot

Applicability

Function Usage

Checks whether the input of each active thread in a warp is not 0.

This API is executed by all active threads in the warp to check the input operand predicate of all active threads and returns a 32-bit unsigned integer. If predicate input by the active threads in the warp is not 0, the bit corresponding to the thread lane ID in the return value is 1. Otherwise, the bit is 0. All active threads in the warp return the same result.

Prototype

1

inline uint32_t asc_ballot(int32_t predicate)

Parameters

Returns

32-bit unsigned integer. If predicate input by the active threads in the warp is not 0, the bit corresponding to the thread lane ID in the return value is 1. Otherwise, the bit is 0.

Restrictions

None

Header File to Be Included

To use this API, the simt_api/device_warp_functions.h header file must be included.

1

#include "simt_api/device_warp_functions.h"

Examples

1
2
3
4
5
6
7
8

__simt_vf__ __launch_bounds__(1024) inline void KernelBallot(__gm__ uint32_t* dst)
{
    // asc_vf_call parameter: dim3{1024, 1, 1}
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int32_t laneId = idx % 32;
    uint32_t result = asc_ballot(laneId); // The return value is 0xfffffffe.
    dst[idx] = result;
}