开发者
下载

简介

在SIMT架构的硬件层面上,处理器以32个线程为一组(一个Warp)来执行线程。如果核函数代码中的条件分支导致Warp内的线程出现发散(Warp Divergence),那么Warp会串行执行每个分支,在执行某个分支时会屏蔽不在该指令路径上的线程。coalesced_group用于获取当前Warp中实际活跃的线程子集。

使用coalesced_group时需关注SIMT架构不支持独立线程调度。

需要包含的头文件

1
#include <simt_api/cooperative_groups.h>

Public成员函数

void sync() const
unsigned long long num_threads() const;
unsigned long long thread_rank() const;
unsigned long long meta_group_size() const;
unsigned long long meta_group_rank() const;
template <typename T>
T shfl(T var, int src_rank) const;
template <typename T>
T shfl_up(T var, unsigned int delta) const;
template <typename T>
T shfl_down(T var, unsigned int delta) const;
int any(int predicate) const;
int all(int predicate) const;
unsigned int ballot(int predicate) const;
unsigned long long size() const;