asc_reduce_add

Applicability

Function Usage

Sums up the val values 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_add(int32_t val)

1

inline uint32_t asc_reduce_add(uint32_t val)

1

inline float asc_reduce_add(float val)

1

inline half asc_reduce_add(half val)

Parameters

Returns

Sum of val input by all threads in the warp.

Restrictions

• If the sum of data overflows, the computing precision of this API is not ensured.
• This API is implemented using the binary algorithm. In some scenarios, the computation result may be different from that of the sequential computation result. Simply put, the computation sequence of (((a + b) + c) + d) is different from that of ((a + b) + (c + d)), which may lead to different final results. This is because in floating-point number computation, each addition operation involves a numerical representation with limited precision. The rounding operation in this process can cause precision loss. Consequently, different addition sequences may lead to different intermediate results, which in turn affects the accuracy of the final computation results.

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 KernelReduceAdd(__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_add(laneId); // The return value is 0+1+2+...+31=496.
     dst[idx] = result;
}