asc_atomic_exch

Applicability

Function Usage

Performs an atomic assignment operation on the address in the Unified Buffer or Global Memory, assigning the specified data to the address in the Unified Buffer or Global Memory.

Prototype

1

inline float asc_atomic_exch(float *address, float val)

1

inline int32_t asc_atomic_exch(int32_t *address, int32_t val)

1

inline uint32_t asc_atomic_exch(uint32_t *address, uint32_t val)

1

inline int64_t asc_atomic_exch(int64_t *address, int64_t val)

1

inline uint64_t asc_atomic_exch(uint64_t *address, uint64_t val)

1

inline half2 asc_atomic_exch(half2 *address, half2 val)

1

inline bfloat16x2_t asc_atomic_exch(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.

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.

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 KernelAtomicExch(__gm__ int32_t* dst, __gm__ int32_t* src)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    asc_atomic_exch(dst + idx, src[idx]);
}