输入关键字搜索
开发者能力认证
鸿蒙高校创新赛
HarmonyOS 7
animation
HarmonyOS行业解决方案
layoutWeight
HarmonyOS模板&组件
bindPopup
bindsheet
CTRL+K
管理中心
我的

HarmonyOS

输入关键字搜索
开发者能力认证
鸿蒙高校创新赛
HarmonyOS 7
animation
HarmonyOS行业解决方案
layoutWeight
HarmonyOS模板&组件
bindPopup
bindsheet
CTRL+K
版本说明

InitBufPool

更新时间: 2026-08-18 12:49
本文导读
展开章节

功能说明

初始化TBufPool内存资源池。本接口适用于内存资源有限时，希望手动指定UB/L1内存资源复用的场景。本接口初始化后在整体内存资源中划分出一块子资源池。划分出的子资源池TBufPool，提供了如下方式进行资源管理：

  • TPipe::InitBufPool的重载接口指定与其他TBufPool子资源池复用。

  • TBufPool::InitBufPool接口对子资源池继续划分。

  • TBufPool::InitBuffer接口分配Buffer。

关于TBufPool的具体介绍及资源划分图示请参考TBufPool。

函数原型

收起
自动换行
深色代码主题
复制
  1. template <class T> 
  2. __aicore__ inline bool InitBufPool(T& bufPool, uint32_t len)
  3. template <class T, class U> 
  4. __aicore__ inline bool InitBufPool(T& bufPool, uint32_t len, U& shareBuf)

参数说明

表1 InitBufPool(T& bufPool, uint32_t len) 原型定义参数说明

展开
参数名称 输入/输出 含义
bufPool 输入 新划分的资源池，类型为TBufPool。
len 输入 新划分资源池长度，单位为Byte，非32Bytes对齐会自动补齐至32Bytes对齐。

表2 InitBufPool(T& bufPool, uint32_t len, U& shareBuf)原型定义参数说明

展开
参数名称 输入/输出 含义
bufPool 输入 新划分的资源池，类型为TBufPool。
len 输入 新划分资源池长度，单位为Byte，非32Bytes对齐会自动补齐至32Bytes对齐。
shareBuf 输入 被复用资源池，类型为TBufPool，新划分资源池与被复用资源池共享起始地址及长度。

支持的型号

Kirin9020系列处理器

Kirin9030系列处理器

KirinX90系列处理器

注意事项

  • 新划分的资源池与被复用资源池的硬件属性需要一致，两者共享起始地址及长度。

  • 输入长度需要小于等于被复用资源池长度。

  • 其他泛用约束参考TBufPool。

返回值

无

调用示例

由于物理内存的大小有限，在计算过程没有数据依赖的场景或数据依赖串行、没有内存踩踏的场景下，可以通过指定内存复用解决资源不足的问题。本示例中Tpipe::InitBufPool初始化子资源池tbufPool1，并且指定tbufPool2复用tbufPool1的起始地址及长度。tbufPool1及tbufPool2的后续计算串行，不存在数据踩踏，实现了内存复用及自动同步的能力。

收起
自动换行
深色代码主题
复制
  1. #include "kernel_operator.h"
  2. class ResetApi {
  3. public:
  4.     __aicore__ inline ResetApi() {}
  5.     __aicore__ inline void Init(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
  6.     {
  7.         src0Global.SetGlobalBuffer((__gm__ half*)src0Gm);
  8.         src1Global.SetGlobalBuffer((__gm__ half*)src1Gm);
  9.         dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
  10.         pipe.InitBufPool(tbufPool1, 196608);
  11.         pipe.InitBufPool(tbufPool2, 196608, tbufPool1);
  12.     }
  13.     __aicore__ inline void Process()
  14.     {
  15.         tbufPool1.InitBuffer(queSrc0, 1, 65536);
  16.         tbufPool1.InitBuffer(queSrc1, 1, 65536);
  17.         tbufPool1.InitBuffer(queDst0, 1, 65536);
  18.         CopyIn();
  19.         Compute();
  20.         CopyOut();
  21.         tbufPool1.Reset();
  22.         tbufPool2.InitBuffer(queSrc2, 1, 65536);
  23.         tbufPool2.InitBuffer(queSrc3, 1, 65536);
  24.         tbufPool2.InitBuffer(queDst1, 1, 65536);
  25.         CopyIn1();
  26.         Compute1();
  27.         CopyOut1();
  28.         tbufPool2.Reset();
  29.     }
  30. private:
  31.     __aicore__ inline void CopyIn()
  32.     {
  33.         AscendC::LocalTensor<half> src0Local = queSrc0.AllocTensor<half>();
  34.         AscendC::LocalTensor<half> src1Local = queSrc1.AllocTensor<half>();
  35.         AscendC::DataCopy(src0Local, src0Global, 512);
  36.         AscendC::DataCopy(src1Local, src1Global, 512);
  37.         queSrc0.EnQue(src0Local);
  38.         queSrc1.EnQue(src1Local);
  39.     }
  40.     __aicore__ inline void Compute()
  41.     {
  42.         AscendC::LocalTensor<half> src0Local = queSrc0.DeQue<half>();
  43.         AscendC::LocalTensor<half> src1Local = queSrc1.DeQue<half>();
  44.         AscendC::LocalTensor<half> dstLocal = queDst0.AllocTensor<half>();
  45.         AscendC::Add(dstLocal, src0Local, src1Local, 512);
  46.         queDst0.EnQue<half>(dstLocal);
  47.         queSrc0.FreeTensor(src0Local);
  48.         queSrc1.FreeTensor(src1Local);
  49.     }
  50.     __aicore__ inline void CopyOut()
  51.     {
  52.         AscendC::LocalTensor<half> dstLocal = queDst0.DeQue<half>();
  53.         AscendC::DataCopy(dstGlobal, dstLocal, 512);
  54.         queDst0.FreeTensor(dstLocal);
  55.     }
  56.     __aicore__ inline void CopyIn1()
  57.     {
  58.         AscendC::LocalTensor<half> src0Local = queSrc2.AllocTensor<half>();
  59.         AscendC::LocalTensor<half> src1Local = queSrc3.AllocTensor<half>();
  60.         AscendC::DataCopy(src0Local, src0Global, 512);
  61.         AscendC::DataCopy(src1Local, src1Global, 512);
  62.         queSrc2.EnQue(src0Local);
  63.         queSrc3.EnQue(src1Local);
  64.     }
  65.     __aicore__ inline void Compute1()
  66.     {
  67.         AscendC::LocalTensor<half> src0Local = queSrc2.DeQue<half>();
  68.         AscendC::LocalTensor<half> src1Local = queSrc3.DeQue<half>();
  69.         AscendC::LocalTensor<half> dstLocal = queDst1.AllocTensor<half>();
  70.         AscendC::Add(dstLocal, src0Local, src1Local, 512);
  71.         queDst1.EnQue<half>(dstLocal);
  72.         queSrc2.FreeTensor(src0Local);
  73.         queSrc3.FreeTensor(src1Local);
  74.     }
  75.     __aicore__ inline void CopyOut1()
  76.     {
  77.         AscendC::LocalTensor<half> dstLocal = queDst1.DeQue<half>();
  78.         AscendC::DataCopy(dstGlobal, dstLocal, 512);
  79.         queDst1.FreeTensor(dstLocal);
  80.     }
  81. private:
  82.     AscendC::TPipe pipe;
  83.     AscendC::TBufPool<AscendC::TPosition::VECCALC> tbufPool1, tbufPool2;
  84.     AscendC::TQue<AscendC::QuePosition::VECIN, 1> queSrc0, queSrc1, queSrc2, queSrc3;
  85.     AscendC::TQue<AscendC::QuePosition::VECOUT, 1> queDst0, queDst1;
  86.     AscendC::GlobalTensor<half> src0Global, src1Global, dstGlobal;
  87. };
  88. extern "C" __global__ __aicore__ void tbufpool_kernel(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
  89. {
  90.     ResetApi op;
  91.     op.Init(src0Gm, src1Gm, dstGm);
  92.     op.Process();
  93. }
相关推荐
文档InitBuffer
文档简介
文档Reset
文档Init
文档TaskPool使用规范
文档TaskPool简介
文档SoundPool (音频池)
文档全局复用：集中化的组件回收与复用
文档@ohos.taskpool (启动任务池)
文档I/O密集型任务开发指导 (TaskPool)
简体中文
华为开发者联盟 版权所有 ©2026
问题反馈
使用条款
华为开发者联盟用户协议
关于华为开发者联盟与隐私的声明
Cookies
开源软件声明