GetTPipePtr

Applicability

Product	Supported/Unsupported
Atlas A3 training products/Atlas A3 inference products	√
Atlas A2 training products/Atlas A2 inference products	√
Atlas 200I/500 A2 inference products	x
Atlas inference product's AI Core	√
Atlas inference product's Vector Core	x
Atlas training products	√

Function Usage

When a TPipe object is created, a globally unique TPipe pointer is set during object initialization. This API is called to obtain the pointer. After obtaining the pointer, you can perform TPipe-related operations.

Prototype

__aicore__ inline AscendC::TPipe* GetTPipePtr()

Restrictions

None

Example

In the following example, a TPipe object is created at the entry to the kernel function. A globally unique TPipe pointer is set during object initialization. When the Init function of the KernelAdd class is called, the TPipe pointer does not need to be explicitly passed. Instead, GetTPipePtr is directly used in the function to obtain the global TPipe pointer for operations such as 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 * 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()
    {
        // Operator kernel logic
        ...
    }
private:
    AscendC::TQue<AscendC::TPosition::VECIN, 2> inQueueX, inQueueY;
    AscendC::TQue<AscendC::TPosition::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();
}

Parent topic: Pipe and Que Frameworks