asc_atomic_add

Applicability

Function Usage

Performs an atomic add operation of the specified data to the data in the Unified Buffer or Global Memory, accumulating the specified data to the data in these memory areas.

Prototype

1

inline int32_t asc_atomic_add(int32_t *address, int32_t val)

1

inline uint32_t asc_atomic_add(uint32_t *address, uint32_t val)

1

inline float asc_atomic_add(float *address, float val)

1

inline int64_t asc_atomic_add(int64_t *address, int64_t val)

1

inline uint64_t asc_atomic_add(uint64_t *address, uint64_t val)

1

inline half asc_atomic_add(half *address, half val)

1

inline bfloat16_t asc_atomic_add(bfloat16_t *address, bfloat16_t val)

1

inline half2 asc_atomic_add(half2 *address, half2 val)

1

inline bfloat16x2_t asc_atomic_add(bfloat16x2_t *address, bfloat16x2_t val)

Parameters

The following describes the memory ranges supported by different data types.

Returns

Initial data on the Unified Buffer or Global Memory.

Note that the return values of the half and bfloat16_t types are not accurate due to underlying hardware restrictions. Do not directly use the return values of these types.

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

To use APIs other than half, half2, bfloat16_t, and bfloat16x2_t, the simt_api/device_atomic_functions.h header file must be included. To use half and half2 APIs, the simt_api/asc_fp16.h header file must be included. To use bfloat16_t and bfloat16x2_t APIs, the simt_api/asc_bf16.h header file must be included.

1

#include "simt_api/device_atomic_functions.h"

1

#include "simt_api/asc_fp16.h"

1

#include "simt_api/asc_bf16.h"

Examples

For SIMD and SIMT programming:

1
2
3
4
5

__simt_vf__ __launch_bounds__(1024) inline void KernelAtomicAdd(__gm__ float* dst, __gm__ float* src)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    asc_atomic_add(dst + idx, src[idx]);
}