同步与内存栅栏简介【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkitSIMT程序中不同Thread的执行进度和内存访问到达顺序可能不同。当线程之间存在数据依赖时需要使用同步接口或内存栅栏接口约束执行顺序和内存可见性。同步机制主要分为两类同步屏障barrier要求同一作用范围内的线程都到达指定位置后程序才能继续执行。内存栅栏memory fence约束调用线程在栅栏前后的内存访问顺序使栅栏前的内存操作按指定范围对其他线程可见。内存栅栏不会等待其他线程到达同一位置。API列表类型接口名作用范围是否阻塞线程功能说明同步屏障asc_syncthreads当前Thread Block是等待当前Thread Block内所有Thread都执行到该同步点并保证同步点前的内存操作对块内线程可见。内存栅栏asc_threadfence_block当前Thread Block否保证调用线程在栅栏前的内存操作按顺序对当前Thread Block内的其他线程可见。内存栅栏asc_threadfence全局范围否保证调用线程在栅栏前的全局内存和共享内存写操作按顺序对其他线程可见。同步屏障asc_syncthreads用于Thread Block内的阶段同步。典型场景是多个Thread先写入Unified Buffer中的共享数据再统一进入下一阶段读取这些数据。__global__ __launch_bounds__(256) void block_reduce(float *out, const float *in) { __ubuf__ float buf[256]; uint32_t tid threadIdx.x; buf[tid] in[blockIdx.x * blockDim.x tid]; // 等待当前Thread Block内所有Thread完成buf写入。 asc_syncthreads(); if (tid 0) { float sum 0.0f; for (uint32_t i 0; i blockDim.x; i) { sum buf[i]; } out[blockIdx.x] sum; } }使用asc_syncthreads时需要注意同一Thread Block内的所有Thread必须都执行到该同步点否则已经到达同步点的Thread会一直等待导致程序无法继续执行。避免在分支中调用asc_syncthreads除非可以保证当前Thread Block内所有Thread都会进入该分支。asc_syncthreads只同步当前Thread Block内的Thread不能用于不同Thread Block之间的全局同步。内存栅栏内存栅栏用于约束调用线程的内存访问顺序。它解决的是“某个线程先写数据、再发布状态或标志”这类可见性问题但不会阻塞其他线程也不会让其他线程自动等待。典型的生产者-消费者场景如下data[idx] value; // 确保data写入先于后续标志位更新对其他线程可见。 asc_threadfence(); asc_atomic_exch(ready, 1U);在上述场景中asc_threadfence保证调用线程在栅栏前的数据写入先于标志位更新对其他线程可见。消费者线程仍需要通过轮询、原子操作或其他同步方式判断ready状态内存栅栏本身不会等待消费者线程。asc_threadfence_block与asc_threadfence的区别在于作用范围不同asc_threadfence_block用于Thread Block内的数据可见性顺序约束适合块内共享数据和Unified Buffer协作场景。asc_threadfence用于更大范围的数据可见性顺序约束适合通过Global Memory发布数据或标志的场景。使用建议需要Thread Block内所有Thread完成某一阶段后再继续执行时使用asc_syncthreads。只需要保证当前Thread的内存写入顺序和可见性时使用asc_threadfence_block或asc_threadfence。不要将内存栅栏当作同步屏障使用。内存栅栏不会等待其他Thread也不保证其他Thread已经执行到同一位置。不要将asc_syncthreads当作跨Thread Block同步接口使用。不同Thread Block之间如需共享状态应结合Global Memory、原子操作和内存栅栏设计同步协议。多个Thread并发写同一地址时asc_syncthreads和内存栅栏都不能替代原子操作需要避免写冲突或使用原子操作保证读改写过程的原子性。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考