__popc

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

Counts the number of bits set to 1 in the binary representation of the input data, starting from the most significant bit to the least significant bit.

Prototype

int __popc(unsigned int x)

int __popc(unsigned long long x)

int __popc(unsigned long x)

Parameters

**Table 1** Parameters
Parameter	Input/Output	Description
x	Input	Source operand.

Returns

The number of bits set to 1 in the binary representation of the input data.

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 KernelPopc(__gm__ int* dst, __gm__ unsigned int* x)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __popc(x[idx]);
}

__simt_vf__ __launch_bounds__(1024) inline void KernelPopc(__gm__ int* dst, __gm__ unsigned long long* x)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __popc(x[idx]);
}

Parent topic: Mathematical Functions