开发者
下载

ballot

产品支持情况

产品

是否支持

Atlas 350 加速卡

Atlas A3 训练系列产品/Atlas A3 推理系列产品

x

Atlas A2 训练系列产品/Atlas A2 推理系列产品

x

Atlas 200I/500 A2 推理产品

x

Atlas 推理系列产品AI Core

x

Atlas 推理系列产品Vector Core

x

Atlas 训练系列产品

x

功能说明

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

函数原型

1
unsigned int ballot(int predicate) const

参数说明

表1 参数说明

参数名

输入/输出

描述

predicate

输入

操作数。

返回值说明

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

约束说明

调用示例

  • SIMT编程场景:
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    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混合编程场景:
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    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
        }
        ...
    }