asc_atomic_max

Applicability

Function Usage

Performs an atomic maximum operation on the data in Unified Buffer or Global Memory, assigning the maximum value between the data in Unified Buffer or Global Memory and the specified data to the address in Unified Buffer or Global Memory.

Prototype

1

inline int32_t asc_atomic_max(int32_t *address, int32_t val)

1

inline uint32_t asc_atomic_max(uint32_t *address, uint32_t val)

1

inline float asc_atomic_max(float *address, float val)

1

inline int64_t asc_atomic_max(int64_t *address, int64_t val)

1

inline uint64_t asc_atomic_max(uint64_t *address, uint64_t val)

1

inline half asc_atomic_max(half *address, half val)

1

inline bfloat16_t asc_atomic_max(bfloat16_t *address, bfloat16_t val)

1

inline half2 asc_atomic_max(half2 *address, half2 val)

1

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

Note that the return values of the half and bfloat16_t types are not accurate due to underlying hardware restrictions. Do not directly use the return values of these types.

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

To use APIs other than half, half2, bfloat16_t, and bfloat16x2_t, the simt_api/device_atomic_functions.h header file must be included. To use half and half2 APIs, the simt_api/asc_fp16.h header file must be included. To use bfloat16_t and bfloat16x2_t APIs, 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 KernelAtomicMax(__gm__ float* dst, __gm__ float* src)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    asc_atomic_max(dst + idx, src[idx]);
}