asc_all

Applicability

Function Usage

Checks whether the inputs of all active threads are not 0.

This API is executed by all active threads in the warp to check the input operand predicate of all active threads. If predicate of all active threads is not 0, 1 is returned. Otherwise, 0 is returned. All active threads in the warp return the same result.

Prototype

1

inline int32_t asc_all(int32_t predicate)

Parameters

Returns

If the inputs of all active threads in the warp are not 0, 1 is returned. Otherwise, 0 is returned.

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 KernelAll(__gm__ int32_t* dst)
{
    // asc_vf_call parameter: dim3{1024, 1, 1}
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int32_t laneId = idx % 32;
    int32_t result = asc_all(laneId); // The return value is 0.
    dst[idx] = result;
}