tiled_partition

产品支持情况

产品	是否支持
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

功能说明

tiled_partition接口用于将一个线程组划分为多个更小、固定大小的子组，以便线程在以更精细的粒度上进行协作。本接口包含带有模板参数和不带模板参数的两类接口，分别用于编译时确定划分大小以及运行时确定划分大小的场景。

函数原型

编译时确定划分大小

template <unsigned int Size, typename ParentT>
thread_block_tile<Size, ParentT> tiled_partition(const ParentT& g)

运行时确定划分大小

thread_group tiled_partition(const thread_group& parent, unsigned int tilesz)

thread_group tiled_partition(const thread_block& parent, unsigned int tilesz)

coalesced_group tiled_partition(const coalesced_group& parent, unsigned int tilesz)

参数说明

表1 编译时确定划分大小的接口参数说明
参数名	输入/输出	描述
Size	输入	模板参数，指定划分出的thread_block_tile组大小。
ParentT	输入	模板参数，被划分的父组g的类型，支持thread_block、thread_block_tile。
g	输入	被划分的父组。

表2 运行时确定划分大小的接口参数说明
参数名	输入/输出	描述
parent	输入	被划分的父组，类型支持thread_block、coalesced_group。
tilesz	输入	指定划分出的子组大小。

返回值说明

返回划分出的子组对象。

约束说明

Size 必须为2^n，且小于32，且小于父组的线程数。

对于编译时确定划分大小的接口，父类中的线程数必须能被Size整除。

调用示例

SIMT编程场景：

using namespace cooperative_groups;
__global__ void simt_kernel(...)
{
    ...
    thread_block block = this_thread_block();
    auto tile4 = tiled_partition<4>(block);
    ...
}

SIMD与SIMT混合编程场景：

using namespace cooperative_groups;
__simt_vf__ inline void simt_kernel(...)
{
    ...
    thread_block block = this_thread_block();
    auto tile4 = tiled_partition<4>(block);
    ...
}

父主题： 协作组