tiled_partition

产品支持情况

• Ascend 950PR/Ascend 950DT：支持

• Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持

• Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持

• Atlas 200I/500 A2 推理产品：不支持

• Atlas 推理系列产品AI Core：不支持
• Atlas 推理系列产品Vector Core：不支持

• Atlas 训练系列产品：不支持

功能说明

tiled_partition API用于将一个线程组划分为多个更小、固定大小的子组，以便线程在以更精细的粒度上进行协作。提供模板和非模板两个版本的接口，分别用于编译时确定划分大小以及运行时确定划分大小的场景。模板版本支持创建跨Warp的thread_block_tile。

函数原型

C++

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

C++

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

C++

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

C++

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

参数说明

表1 模板版本参数说明

参数名	输入/输出	描述
g	输入	被划分的父组，类型只能是thread_block或thread_block_tile。
Size	输入	模板参数，指定划分出的thread_block_tile组大小。

表2 非模板版本参数说明

参数名	输入/输出	描述
parent	输入	被划分的父组，类型只能是thread_block或coalesced_group。
tilesz	输入	指定划分出的子组大小。

返回值说明

返回划分出的子组对象。

约束说明

• 模板版本中，Size必须是，当前可选值范围：1、2、4、8、16、32、64、128、256、512、1024、2048。
• 非模板版本仅支持创建Size <= 32的子组。
• 当要创建跨Warp的thread_block_tile时，用户需创建block_tile_memory作为临时存储，并通过带scratch参数的this_thread_block创建父thread_block。传入的block_tile_memory对象必须位于Global Memory或Unified Buffer，不能传入栈空间中创建的对象。使用位于Unified Buffer的对象性能优于位于Global Memory的。
• 对于模板版本的接口，父组中的线程数必须能被Size整除。并且如果父组是thread_block_tile，则Size必须小于父组大小。

调用示例

• SIMT编程场景：

  C++

  using namespace cooperative_groups;
  __global__ void simt_kernel(...)
  {
      ...
      // 无需使用跨warp的thread_block_tile
      thread_block block = this_thread_block();
      auto tile4 = tiled_partition<4>(block);
  
      // 需要使用跨warp的thread_block_tile
      __ubuf__ block_tile_memory<1024> scratch;
      thread_block block_with_memory = this_thread_block(scratch);
      auto tile64 = tiled_partition<64>(block_with_memory);
      ...
  }
  

• SIMD与SIMT混合编程场景：

  C++

  using namespace cooperative_groups;
  __simt_vf__ inline void simt_kernel(...)
  {
      ...
      // 无需使用跨warp的thread_block_tile
      thread_block block = this_thread_block();
      auto tile4 = tiled_partition<4>(block);
  
      // 需要使用跨warp的thread_block_tile
      __ubuf__ block_tile_memory<1024> scratch;
      thread_block block_with_memory = this_thread_block(scratch);
      auto tile64 = tiled_partition<64>(block_with_memory);
      ...
  }