asc_shfl

Applicability

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

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

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;
}