asc_threadfence

Applicability

Function Usage

In the SIMT programming paradigm, different threads may read and write the same memory, which may cause data races. The asc_threadfence API is used to ensure the time sequence of write operations when different threads access the same global and shared memory. It does not block threads and only ensures the visibility sequence of memory operations.

Prototype

1

inline void asc_threadfence()

Parameters

None

Returns

None

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

To use this API, the simt_api/device_sync_functions.h header file must be included.

1

#include "simt_api/device_sync_functions.h"

Examples

1
2
3
4
5
6

__simt_vf__ __launch_bounds__(1024) inline void KernelThreadFence(__gm__ float* dst, __gm__ float* src)
{
    src[0] = src[0] + 1;
    asc_threadfence(); // asc_threadfence() ensures that the write operation sequence of the current thread is globally visible.
    dst[0] = src[0];
}