asc_threadfence

产品支持情况

功能说明

在SIMT编程范式中，来自不同线程对同一份内存的读写操作可能造成数据竞争（Data Race），这是由于NPU架构的特性，线程的执行顺序和内存访问顺序可能不一致，从而导致其他线程看到的数据更新顺序与实际写入顺序不同。

内存栅栏类接口可以解决这类多线程同步读写共享内存的问题，其中asc_threadfence接口用于保证所有线程对同一份全局、共享内存的访问过程中，写入操作的时序性和可见性。

本接口一般适用的场景包括：

• 生产者-消费者同步场景：确保数据写入在标志位设置之前完成。
• 原子计数器场景：在原子计数器更新前确保相关数据已写入。
• 指针更新场景：链表、树等数据结构操作时确保数据一致性
• 多线程共享数据：确保共享数据的更新顺序对其他线程可见。

下图展示了没有内存栅栏时可能出现的内存可见性问题：

通过使用asc_threadfence()接口，可解决上述内存可见性问题：

函数原型

1

inline void asc_threadfence()

参数说明

无

返回值说明

无

约束说明

无

需要包含的头文件

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

1

#include "simt_api/device_sync_functions.h"

调用示例

• SIMT编程场景：

  1 2 3 4 5 6

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

• SIMD与SIMT混合编程场景：

  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()保证本线程的写操作顺序对全局可见 dst[0] = src[0]; }