CANN/asc-devkit SIMT内存栅栏接口文档

asc_threadfence【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品不支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明在SIMT编程范式中来自不同线程对同一份内存的读写操作可能造成数据竞争Data Race这是由于NPU架构的特性线程的执行顺序和内存访问顺序可能不一致从而导致其他线程看到的数据更新顺序与实际写入顺序不同。内存栅栏类接口可以解决这类多线程同步读写共享内存的问题其中asc_threadfence接口用于保证所有线程对同一份全局、共享内存的访问过程中写入操作的时序性和可见性。关键特征该接口不会阻塞线程仅保证内存操作的可见性顺序。作用范围是全局的调用后确保调用线程在asc_threadfence()之前的所有全局内存和共享内存写操作对其他线程可见。本接口一般适用的场景包括生产者-消费者同步场景确保数据写入在标志位设置之前完成。原子计数器场景在原子计数器更新前确保相关数据已写入。指针更新场景链表、树等数据结构操作时确保数据一致性多线程共享数据确保共享数据的更新顺序对其他线程可见。下图展示了没有内存栅栏时可能出现的内存可见性问题通过使用asc_threadfence()接口可解决上述内存可见性问题函数原型inline void asc_threadfence()参数说明无返回值说明无约束说明无需要包含的头文件使用该接口需要包含simt_api/device_sync_functions.h头文件。#include simt_api/device_sync_functions.h调用示例完整样例请参考MemoryFence样例。SIMT编程场景__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混合编程场景__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]; }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考