__ffs

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

Searches for the position of the first bit whose value is 1 in the binary representation of the input data starting from the least significant bit and returns the index of the position. The index starts from 1. If there is no 1 in the binary data, 0 is returned.

Prototype

1
int __ffs(int x)
1
int __ffs(long long x)
1
int __ffs(long x)

Parameters

Table 1 Parameters

Parameter

Input/Output

Description

x

Input

Source operand.

Returns

The position of the first bit whose value is 1 in the binary representation of the input data starting from the least significant bit.

  • If x is 0, the return value is 0.
  • If x is 1, the return value is 1.

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:
1
2
3
4
5
__simt_vf__ __launch_bounds__(1024) inline void KernelFfs(__gm__ int* dst, __gm__ int* x)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __ffs(x[idx]);
}
1
2
3
4
5
__simt_vf__ __launch_bounds__(1024) inline void KernelFfs(__gm__ int* dst, __gm__ long long* x)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __ffs(x[idx]);
}