asc_atomic_cas

Applicability

Function Usage

Performs an atomic compare-and-swap operation on the address value in the Unified Buffer or Global Memory. If the address value is equal to the specified value compare, the address value is set to the specified value val. Otherwise, the address value remains unchanged.

Prototype

1

inline float asc_atomic_cas(float *address, float compare, float val)

1

inline int32_t asc_atomic_cas(int32_t *address, int32_t compare, int32_t val)

1

inline uint32_t asc_atomic_cas(uint32_t *address, uint32_t compare, uint32_t val)

1

inline int64_t asc_atomic_cas(int64_t *address, int64_t compare, int64_t val)

1

inline uint64_t asc_atomic_cas(uint64_t *address, uint64_t compare, uint64_t val)

1

inline half2 asc_atomic_cas(half2 *address, half2 compare, half2 val)

1

inline bfloat16x2_t asc_atomic_cas(bfloat16x2_t *address, bfloat16x2_t compare, 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 KernelAtomicCas(__gm__ int32_t* dst, __gm__ int32_t* src0, __gm__ int32_t* src1)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    asc_atomic_cas(dst + idx, src0[idx], src1[idx]);
}