asc_shfl

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

Obtains the value of var input by the specified thread src_lane in the warp for swapping. If the target thread is inactive, the uninitialized value in the register is obtained. The width parameter is used to group threads in the warp. The width parameter specifies the group width of the 32 threads involved in the swapping. The default value is 32, indicating that all threads are divided into one group.

In the multi-group scenario (the value of width is less than 32), the swapping operation of each group is independent. The current thread obtains the value of var of the thread whose ID is src_lane in the group. If the value of src_lane is less than that of width, src_lane is the logical lane ID relative to the position of the start thread in the group. If the value of src_lane is greater than or equal to that of width, the result of src_lane%width is the logical lane ID relative to the position of the start thread in the group.

For example, if 32 active threads in a warp call the asc_shfl(LaneId, 1, 16) API, the return value of each thread is the value of var of the thread whose ID is 1 in the group where the current thread is located.

Figure 1 asc_shfl result

Prototype

1
inline int32_t asc_shfl(int32_t var, int32_t src_lane, int32_t width = warpSize)
1
inline uint32_t asc_shfl(uint32_t var, int32_t src_lane, int32_t width = warpSize)
1
inline float asc_shfl(float var, int32_t src_lane, int32_t width = warpSize)
1
inline int64_t asc_shfl(int64_t var, int32_t src_lane, int32_t width = warpSize)
1
inline uint64_t asc_shfl(uint64_t var, int32_t src_lane, int32_t width = warpSize)
1
inline half asc_shfl(half var, int32_t src_lane, int32_t width = warpSize)
1
inline half2 asc_shfl(half2 var, int32_t src_lane, int32_t width = warpSize)

Parameters

Table 1 Parameters

Parameter

Input/Output

Description

var

Input

Input operand used by a thread for swapping.

src_lane

Input

Lane ID of the thread whose var value is expected to be retrieved. The value range is [0, 32).

width

Input

Group width of the thread involved in the swapping within the warp. The default value is 32. The value range of width is (0, 32]. It must be a multiple of 2.

Returns

  • Value of var of the specified thread in the warp.
  • Uninitialized undefined value.

Restrictions

None

Header File to Be Included

To use APIs other than half and half2, the simt_api/device_warp_functions.h header file must be included. To use half and half2 APIs, the simt_api/asc_fp16.h header file must be included.

1
#include "simt_api/device_warp_functions.h" 
1
#include "simt_api/asc_fp16.h"

Examples

For SIMD and SIMT programming:
1
2
3
4
5
6
7
8
9
__simt_vf__ __launch_bounds__(1024) void KernelShfl(__gm__ int32_t* dst)
{
    // asc_vf_call parameter: dim3{1024, 1, 1}
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int32_t laneId = idx % 32;
    // The return value of threads 0–15 is 1, and that of threads 16–31 is 17.
    int32_t result = asc_shfl(laneId, 1, 16);
    dst[idx] = result;
}