asc_shfl_down

Applicability

Function Usage

Obtains the var value input by the thread that is delta lanes backward from the current thread in the warp (current LaneId + delta) 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. Each thread obtains the value of var of the thread with the backward offset delta relative to the current thread in the group. If the ID of the thread with the backward offset delta relative to the current thread, that is, lane ID+delta, is greater than the maximum lane ID of the group, the value of var of the current thread is returned.

For example, if 32 active threads in a warp call the asc_shfl_down(LaneId, 2, 16) API, the return value of each thread is the value of var of the thread with the ID (lane ID of the current thread+2), or the value of var of the current thread.

Figure 1 asc_shfl_down result

Prototype

1

inline int32_t asc_shfl_down(int32_t var, uint32_t delta, int32_t width = warpSize)

1

inline uint32_t asc_shfl_down(uint32_t var, uint32_t delta, int32_t width = warpSize)

1

inline float asc_shfl_down(float var, uint32_t delta, int32_t width = warpSize)

1

inline int64_t asc_shfl_down(int64_t var, uint32_t delta, int32_t width = warpSize)

1

inline uint64_t asc_shfl_down(uint64_t var, uint32_t delta, int32_t width = warpSize)

1

inline half asc_shfl_down(half var, uint32_t delta, int32_t width = warpSize)

1

inline half2 asc_shfl_down(half2 var, uint32_t delta, 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
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

__simt_vf__ __launch_bounds__(1024) inline void KernelShflDown(__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 values of threads 0 to 15 are {2,3,4,5,6,7,8,9,10,11,12,13,14,15,14,15}.
    // The return values of threads 16 to 31 are {18,19,20,21,22,23,24,25,26,27,28,29,30,31,30,31}.
    int32_t result = asc_shfl_down(laneId, 2, 16);
    dst[idx] = result;
}

// asc_shfl_down implements reducesum.
__simt_vf__ __launch_bounds__(1024) inline void KernelShflDownReduceSum(__gm__ int32_t* dst)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int32_t laneId = idx % 32;
    int32_t value = laneId;

    value += asc_shfl_down(value, 1, 31); // 1
    value += asc_shfl_down(value, 2, 31); // 2
    value += asc_shfl_down(value, 4, 31); // 4
    value += asc_shfl_down(value, 8, 31); // 8
    value += asc_shfl_down(value, 16, 31); // 16

    dst[idx] = value;
}