__byte_perm

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

Combines two 4-byte uint32_t inputs into an 8-byte 64-bit integer. The selector s specifies which 4 bytes to select, and these 4 bytes are concatenated into a uint32_t integer from the least significant bit to the most significant bit. The implementation logic is as follows:

// The following is the BytePerm(x, y, s) calculation logic represented in C++.
uint64_t tmp64 = ((uint64_t)y << 32) | x; // x,y concatenated into a uint64 integer

uint8_t selector0 = (s >> 0) & 0x7; // The value range of selector0 is [0, 7].
uint8_t selector1 = (s >> 4) & 0x7;
uint8_t selector2 = (s >> 8) & 0x7;
uint8_t selector3 = (s >> 12) & 0x7;

uint8_t byte0 = (tmp64 >> (selector0 * 8)) & 0xFF; // Select the selector0 byte in tmp64.
uint8_t byte1 = (tmp64 >> (selector1 * 8)) & 0xFF;
uint8_t byte2 = (tmp64 >> (selector2 * 8)) & 0xFF;
uint8_t byte3 = (tmp64 >> (selector3 * 8)) & 0xFF;

// result is the return value of BytePerm. The result is obtained by concatenating the corresponding bytes in sequence.
uint32_t result = byte0 | (byte1 << 8) | (byte2 << 16) | (byte3 << 24);

Prototype

unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)

Parameters

**Table 1** Parameters
Parameter	Input/Output	Description
x	Input	Source operand of the uint32_t type. It is combined with y to form a 64-bit integer. The [0:31] bit of the integer is x.
y	Input	Source operand of the uint32_t type. It is combined with x to form a 64-bit integer. The [32:63] bit of the integer is y.
s	Input	Selector of the uint32_t type, which specifies how to extract 4-byte data from the 64-bit integer consisting of 8 bytes formed by x and y. Specifically, the data represented by s[0:3], s[4:7], s[8:11], and s[12:15] specifies the index values 0 to 7 of the selected bytes in the 8-byte integer.

Returns

Integer of the uint32_t type selected by the selector s.

If x is 0, y is 0, and s is 0, the return value is 0.
If x is 1, y is 1, and s is 1, the return value is 16843008.

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

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

Examples

For SIMD and SIMT programming:

__simt_vf__ __launch_bounds__(1024) inline void KernelByte_perm(__gm__ unsigned int* dst, __gm__ unsigned int* x, __gm__ unsigned int* y, __gm__ unsigned int* s)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __byte_perm(x[idx], y[idx], s[idx]);
}

Parent topic: Mathematical Functions