asc_shfl_xor

Applicability

Function Usage

Obtains the var value input by the thread corresponding to dstLaneId (obtained by performing an XOR operation on the current LaneId and the input lane_mask: LaneId^lane_mask) 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 (width < 32), each thread obtains the value of var of the thread corresponding to dstLaneId in the current group or a group with a smaller thread ID. That is, if dstLaneId is less than the start lane ID of the group where the current thread is located, the thread corresponding to dstLaneId is in a group with a smaller thread ID. In this case, the value of var of the thread corresponding to dstLaneId can be obtained. If dstLaneId is greater than the maximum lane ID of the group where the current thread is located, the value of var of the current thread is returned.

For example, if 32 active threads in a warp call the asc_shfl_xor(LaneId, 1, 16) API, the return value of each thread is the var value of the thread corresponding to the lane ID of the current thread^1.

Figure 1 asc_shfl_xor result

Prototype

1

inline int32_t asc_shfl_xor(int32_t var, int32_t lane_mask, int32_t width = warpSize)

1

inline uint32_t asc_shfl_xor(uint32_t var, int32_t lane_mask, int32_t width = warpSize)

1

inline float asc_shfl_xor(float var, int32_t lane_mask, int32_t width = warpSize)

1

inline int64_t asc_shfl_xor(int64_t var, int32_t lane_mask, int32_t width = warpSize)

1

inline uint64_t asc_shfl_xor(uint64_t var, int32_t lane_mask, int32_t width = warpSize)

1

inline half asc_shfl_xor(half var, int32_t lane_mask, int32_t width = warpSize)

1

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

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

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

    dst[idx] = value;
}