asc_threadfence_block

产品支持情况

功能说明

在同一Thread Block内，多个线程共享Unified Buffer（256KB共享内存），线程间也会存在数据竞争问题。

与asc_threadfence接口类似，asc_threadfence_block函数用于保证当前线程的内存读写操作在同一Thread Block范围内的可见性顺序，确保某一线程在调用asc_threadfence_block()之前的所有内存读写操作对同一线程块内的其他线程可见。该函数不会阻塞当前线程的执行，仅建立Block范围内的内存操作可见性顺序约束。asc_threadfence_block提供Block范围内的可见性保证，其实现仅需要确保Block内共享的缓存层级一致性，无需进行全局内存屏障操作。

函数原型

1

inline void asc_threadfence_block()

参数说明

无

返回值说明

无

约束说明

无

需要包含的头文件

使用该接口需要包含"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 KernelThreadFenceBlock(float* dst, float* src) { src[0] = src[0] + 1; asc_threadfence_block(); dst[0] = src[0]; }

• SIMD与SIMT混合编程场景：

  1 2 3 4 5 6

  __simt_vf__ __launch_bounds__(1024) inline void KernelThreadFenceBlock(__gm__ float* dst, __gm__ float* src) { src[0] = src[0] + 1; asc_threadfence_block(); dst[0] = src[0]; }