asc_atomic_sub

Applicability

Product	Supported
Atlas 350 Accelerator Card	√
Atlas A3 training product/Atlas A3 inference product	x
Atlas A2 training product/Atlas A2 inference product	x
Atlas 200I/500 A2 inference product	x
Atlas inference product AI Core	x
Atlas inference product Vector Core	x
Atlas training product	x

Function Usage

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

Prototype

inline int32_t asc_atomic_sub(int32_t *address, int32_t val)

inline uint32_t asc_atomic_sub(uint32_t *address, uint32_t val)

inline float asc_atomic_sub(float *address, float val)

inline int64_t asc_atomic_sub(int64_t *address, int64_t val)

inline uint64_t asc_atomic_sub(uint64_t *address, uint64_t val)

inline half2 asc_atomic_sub(half2 *address, half2 val)

inline bfloat16x2_t asc_atomic_sub(bfloat16x2_t *address, bfloat16x2_t val)

Parameters

**Table 1** Parameters
Parameter	Input/Output	Description
address	Output	Address of the Unified Buffer or Global Memory.
val	Input	Source operand.

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

**Table 2** Memory ranges supported by different data types
Parameter Data Type	Supported Memory Space
int32_t, uint32_t, float, half2, and bfloat16x2_t	Unified Buffer and Global Memory
int64_t and uint64_t	Global Memory

Returns

Initial data on the Unified Buffer or Global Memory.

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

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

#include "simt_api/device_atomic_functions.h"

#include "simt_api/asc_fp16.h"

#include "simt_api/asc_bf16.h"

Examples

For SIMD and SIMT programming:

In SIMD and SIMT programming scenarios, the address space qualifier must be explicitly used to indicate the address space. __gm__ indicates the Global Memory memory space, and __ubuf__ indicates the Unified Buffer memory space.

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

Parent topic: Atomic Function