CANN/asc-devkit __popc函数文档
2026/7/18 12:17:26 网站建设 项目流程

__popc

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言,原生支持C和C++标准规范,主要由类库和语言扩展层构成,提供多层级API,满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

产品支持情况

  • Ascend 950PR/Ascend 950DT:支持
  • Atlas A3 训练系列产品/Atlas A3 推理系列产品:不支持
  • Atlas A2 训练系列产品/Atlas A2 推理系列产品:不支持
  • Atlas 200I/500 A2 推理产品:不支持
  • Atlas 推理系列产品AI Core:不支持
  • Atlas 推理系列产品Vector Core:不支持
  • Atlas 训练系列产品:不支持

功能说明

统计输入数据从二进制的高位到低位比特位为1的数量。

函数原型

int __popc(unsigned int x)
int __popc(unsigned long long x)
int __popc(unsigned long x)

参数说明

表1参数说明

参数名输入/输出描述
x输入源操作数。

返回值说明

输入数据二进制中比特位为1的数量。

约束说明

需要包含的头文件

使用该接口需要包含"simt_api/device_functions.h"头文件。

#include "simt_api/device_functions.h"

调用示例

  • SIMT编程场景:

    __global__ __launch_bounds__(1024) void KernelPopc(int* dst, unsigned int* x) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = __popc(x[idx]); }
    __global__ __launch_bounds__(1024) void KernelPopc(int* dst, unsigned long long* x) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = __popc(x[idx]); }
  • SIMD与SIMT混合编程场景:

    __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]); }

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言,原生支持C和C++标准规范,主要由类库和语言扩展层构成,提供多层级API,满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询