GetTPipePtr
函数原型
1 | __aicore__ inline AscendC::TPipe* GetTPipePtr() |
支持的型号
Atlas 训练系列产品
Atlas推理系列产品AI Core
Atlas A2训练系列产品/Atlas 800I A2推理产品
注意事项
无
调用示例
如下样例中,在核函数入口处创建TPipe对象,对象初始化会设置全局唯一的TPipe指针。在调用KernelAdd类Init函数时,无需显示传入TPipe指针,而是在函数内直接使用GetTPipePtr获取全局TPipe指针,用来做InitBuffer等操作。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 | class KernelAdd { public: __aicore__ inline KernelAdd() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { xGm.SetGlobalBuffer((__gm__ half *)x + 2048 * AscendC::GetBlockIdx(), 2048); yGm.SetGlobalBuffer((__gm__ half *)y + 2048 * AscendC::GetBlockIdx(), 2048); zGm.SetGlobalBuffer((__gm__ half *)z + 2048 * AscendC::GetBlockIdx(), 2048); GetTPipePtr()->InitBuffer(inQueueX, 2, 128 * sizeof(half)); GetTPipePtr()->InitBuffer(inQueueY, 2, 128 * sizeof(half)); GetTPipePtr()->InitBuffer(outQueueZ, 2, 128 * sizeof(half)); } __aicore__ inline void Process() { // 算子kernel逻辑 ... } private: AscendC::TQue<AscendC::QuePosition::VECIN, 2> inQueueX, inQueueY; AscendC::TQue<AscendC::QuePosition::VECOUT, 2> outQueueZ; AscendC::GlobalTensor<half> xGm, yGm, zGm; }; extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { AscendC::TPipe pipe; KernelAdd op; op.Init(x, y, z); op.Process(); } |
父主题: 内存管理与同步控制