asc_syncthreads

产品支持情况

功能说明

等待当前线程块内所有线程代码都执行到该函数位置。

函数原型

1

inline void asc_syncthreads()

参数说明

无

返回值说明

无

约束说明

无

需要包含的头文件

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

1

#include "simt_api/device_sync_functions.h"

调用示例

• SIMT编程场景：

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17

  __global__ __launch_bounds__(1024) void KernelSyncThreads(float* dst, int count) { int idx = threadIdx.x; if (idx > 0 && idx < count) { dst[idx] = 1; } // 等待block内所有thread都执行到当前代码 asc_syncthreads(); if (idx == 0) { dst[0] = 0; for(int i = 1023; i > 0; i--) { dst[0] += dst[i]; } } }

  1 2	输出结果: [1023, 1, 1, 1 …]

• SIMD与SIMT混合编程场景：

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17

  __simt_vf__ __launch_bounds__(1024) inline void KernelSyncThreads(__gm__ float* dst, int count) { int idx = threadIdx.x; if (idx > 0 && idx < count) { dst[idx] = 1; } // 等待block内所有thread都执行到当前代码 asc_syncthreads(); if (idx == 0) { dst[0] = 0; for(int i = 1023; i > 0; i--) { dst[0] += dst[i]; } } }

  1 2	输出结果: [1023, 1, 1, 1 …]