asc_atomic_dec

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 decrement operation on the address value in the Unified Buffer or Global Memory. If the address value is 0 or greater than the specified value val, the address value is set to val. Otherwise, the address value is decremented by 1.

Prototype

inline uint32_t asc_atomic_dec(uint32_t *address, uint32_t val)

inline uint64_t asc_atomic_dec(uint64_t *address, uint64_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
uint32_t	Unified Buffer and Global Memory
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 this API, the simt_api/device_atomic_functions.h header file must be included.

#include "simt_api/device_atomic_functions.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 KernelAtomicDec(__gm__ uint32_t* dst, __gm__ uint32_t* src)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    asc_atomic_dec(dst + idx, src[idx]);
}

Parent topic: Atomic Function