asc_threadfence

产品支持情况

产品	是否支持
Ascend 950PR/Ascend 950DT	√
Atlas A3 训练系列产品/Atlas A3 推理系列产品	x
Atlas A2 训练系列产品/Atlas A2 推理系列产品	x

功能说明

在SIMT编程范式中，来自不同线程对同一份内存的读写操作可能造成数据竞争（Data Race）。asc_threadfence接口用于保证不同线程对同一份全局、共享内存的访问过程中，写入操作的时序性。它不会阻塞线程，仅保证内存操作的可见性顺序。

函数原型

__simt_callee__ inline void asc_threadfence()

参数说明

无

返回值说明

无

约束说明

无

需要包含的头文件

使用该接口需要包含"simt_api/device_sync_functions.h"头文件。

#include "simt_api/device_sync_functions.h"

调用示例

__simt_vf__ __launch_bounds__(1024) inline void KernelThreadFence(__gm__ float* dst, __gm__ float* src)
{
    src[0] = src[0] + 1;
    asc_threadfence(); // asc_threadfence()保证本线程的写操作顺序对全局可见
    dst[0] = src[0];
}