asc_reduce_min

Applicability

Function Usage

Computes the minimum value of val input by all active threads in the warp. All active threads in the warp return the same result.

Prototype

1

inline int32_t asc_reduce_min(int32_t val)

1

inline uint32_t asc_reduce_min(uint32_t val)

1

inline float asc_reduce_min(float val)

1

inline half asc_reduce_min(half val)

Parameters

Returns

Minimum value of val input by all active threads in the warp.

Restrictions

None

Header File to Be Included

To use APIs other than half, the simt_api/device_warp_functions.h header file must be included. To use the half API, the simt_api/asc_fp16.h header file must be included.

1

#include "simt_api/device_warp_functions.h" 

1

#include "simt_api/asc_fp16.h"

Examples

1
2
3
4
5
6
7
8

__simt_vf__ __launch_bounds__(1024) inline void KernelReduceMin(__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_reduce_min(laneId); // The return value is 0.
    dst[idx] = result;
}