asc_atomic_and

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 AND (&) operation on the address value in the Unified Buffer or Global Memory and the specified value val, assigning the result to the Unified Buffer or Global Memory.

Prototype

inline int32_t asc_atomic_and(int32_t *address, int32_t val)

inline uint32_t asc_atomic_and(uint32_t *address, uint32_t val)

inline int64_t asc_atomic_and(int64_t *address, int64_t val)

inline uint64_t asc_atomic_and(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
int32_t and uint32_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 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 KernelAtomicAnd(__gm__ int32_t* dst, __gm__ int32_t* src)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    asc_atomic_and(dst + idx, src[idx]);
}

Parent topic: Atomic Function