asc_syncthreads

Applicability

Product	Supported
Atlas 350 Accelerator Card	√
Atlas A3 training product/Atlas A3 inference product	x
Atlas A2 training product/Atlas A2 inference product	x
Atlas 200I/500 A2 inference product	x
Atlas inference product AI Core	x
Atlas inference product Vector Core	x
Atlas training product	x

Function Usage

Waits until all threads in the current thread block have reached this function.

Prototype

inline void asc_syncthreads()

Parameters

None

Returns

None

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

To use this API, the simt_api/device_sync_functions.h header file must be included.

#include "simt_api/device_sync_functions.h"

Examples

For SIMD and SIMT programming:

__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;
     }
     
     // Wait until all threads in the block have reached the current code.
     asc_syncthreads();
        
     if (idx == 0) {
         dst[0] = 0;
         for(int i = 1023; i > 0; i--) {
             dst[0] += dst[i];
         }
     }
}

Output:
[1023, 1, 1, 1 ...]

Parent topic: Synchronization Functions