thread_block_tile构造函数

产品支持情况

• Ascend 950PR/Ascend 950DT：支持

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

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

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

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

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

功能说明

thread_block_tile不提供默认构造函数，可通过tiled_partition接口从另一个协作组中划分得到。

函数原型

C++

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

参数说明

表1 参数说明

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

返回值说明

返回划分后当前线程所属的thread_block_tile组。

约束说明

• Size必须是，当前可选值范围：1、2、4、8、16、32、64、128、256、512、1024、2048。
• 当Size大于32时，即Size为64、128、256、512、1024、2048的跨Warp协作场景，必须使用block_tile_memory创建父thread_block。并且用于创建父thread_block的block_tile_memory对象必须位于Global Memory或Unified Buffer，不能是在栈空间中创建的对象。使用位于Unified Buffer的对象性能优于位于Global Memory的。
• 对于模板版本的接口，父组中的线程数必须能被Size整除。并且Size必须小于父组大小。

调用示例

• SIMT编程场景：

  C++

  using namespace cooperative_groups;
  constexpr int THREAD_NUM = 1024;
  __global__ void simt_kernel(...)
  {
      ...
      // 创建Size<=32的thread_block_tile
      thread_block block = this_thread_block();
      thread_block_tile<32> tile32 = tiled_partition<32>(block);              // 按照32个线程为一组划分thread_block
      auto tile32_auto = tiled_partition<32>(block);                          // 建议使用auto管理返回对象
      thread_block_tile<4, thread_block> tile4 = tiled_partition<4>(block);   // 按照4个线程为一组划分thread_block，对象类型中保留父组信息
  
      // 创建Size>32的thread_block_tile
      __ubuf__ block_tile_memory<THREAD_NUM> scratch;
      thread_block block_with_memory = this_thread_block(scratch);
      auto tile64 = tiled_partition<64>(block_with_memory);                   // 按照64个线程为一组划分thread_block
      ...
  }
  
  int main()
  {
      ...
      simt_kernel<<<dim3(1), dim3(THREAD_NUM), 0, stream>>>(...);             // 线程块中创建的线程数必须小于等于 block_tile_memory 的模板参数
      ...
  }
  

• SIMD与SIMT混合编程场景：

  C++

  constexpr int THREAD_NUM = 1024;
  using namespace cooperative_groups;
  __simt_vf__ inline void simt_kernel(...)
  {
      ...
      // 创建Size<=32的thread_block_tile
      thread_block block = this_thread_block();
      thread_block_tile<32> tile32 = tiled_partition<32>(block);              // 按照32个线程为一组划分thread_block
      auto tile32_auto = tiled_partition<32>(block);                          // 建议使用auto管理返回对象
      thread_block_tile<4, thread_block> tile4 = tiled_partition<4>(block);   // 按照4个线程为一组划分thread_block，对象类型中保留父组信息
  
      // 创建Size>32的thread_block_tile
      __ubuf__ block_tile_memory<THREAD_NUM> scratch;
      thread_block block_with_memory = this_thread_block(scratch);
      auto tile64 = tiled_partition<64>(block_with_memory);                   // 使用UB临时存储划分thread_block
      ...
  }
  
  __global__ __vector__ void global_kernel(...)
  {
      ...
      asc_vf_call<simt_kernel>(dim3(THREAD_NUM), ...);                        // 拉起SIMT VF时配置的线程数必须小于等于 block_tile_memory 的模板参数
      ...
  }