asc_activemask

Applicability

Function Usage

Checks whether all threads in the warp are active.

Returns a 32-bit unsigned integer. If a thread in the warp is active (a finished thread is inactive), 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_activemask()

Parameters

None

Returns

32-bit unsigned integer. If a thread in the warp is active, 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

__simt_vf__ __launch_bounds__(1024) inline void KernelActiveMask(__gm__ uint32_t* dst)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    // asc_vf_call parameter: dim3{1024, 1, 1}
    uint32_t result = asc_activemask(); // The return value is 0xffffffff.
    dst[idx] = result;
}