GetTPipePtr
函数原型
__aicore__ inline AscendC::TPipe* GetTPipePtr();
支持的型号
Atlas 训练系列产品
Atlas推理系列产品AI Core
Atlas A2训练系列产品/Atlas 800I A2推理产品
注意事项
无
调用示例
如下样例中,在核函数入口处创建TPipe对象,对象初始化会设置全局唯一的TPipe指针。在调用KernelAdd类Init函数时,无需显示传入TPipe指针,而是在函数内直接使用GetTPipePtr获取全局TPipe指针,用来做InitBuffer等操作。
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 * GetBlockIdx(), 2048); yGm.SetGlobalBuffer((__gm__ half *)y + 2048 * GetBlockIdx(), 2048); zGm.SetGlobalBuffer((__gm__ half *)z + 2048 * 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: TQue<QuePosition::VECIN, 2> inQueueX, inQueueY; TQue<QuePosition::VECOUT, 2> outQueueZ; GlobalTensor<half> xGm, yGm, zGm; }; extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { TPipe pipe; KernelAdd op; op.Init(x, y, z); op.Process(); }
父主题: 内存管理与同步控制