Skip to content

ballot

产品支持情况

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

功能说明

判断coalesced_group组内每个线程的输入是否非零。

函数原型

C++
unsigned int ballot(int predicate) const

参数说明

表1 参数说明

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

返回值说明

32-bit的无符号整数:若coalesced_group组内线程输入的predicate不为0,则返回值中与线程rank对应的bit位为1,否则为0。

约束说明

调用示例

  • SIMT编程场景:

    C++
    using namespace cooperative_groups;
    __global__ void simt_kernel(...)
    {
        ...
        if (threadIdx.x % 2 == 0) {
            coalesced_group active = coalesced_threads();
            uint32_t result = active.ballot(1); // 返回0xFFFF
        }
        ...
    }
    
  • SIMD与SIMT混合编程场景:

    C++
    using namespace cooperative_groups;
    __simt_vf__ inline void simt_kernel(...)
    {
        ...
        if (threadIdx.x % 2 == 0) {
            coalesced_group active = coalesced_threads();
            uint32_t result = active.ballot(1); // 返回0xFFFF
        }
        ...
    }
    

免责声明:本站内容由 asc-devkit 仓 master 分支自动编译生成,属于持续开发版本,可能存在缺陷,仅供预览与参考。如需稳定及商用资料,请查阅官方 昇腾社区