文档

GetTPipePtr

Function Usage

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

Prototype

1
__aicore__ inline AscendC::TPipe* GetTPipePtr()

Availability

Atlas Training Series Product

Precautions

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.

 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()
    {
        // Operator kernel logic
        ...
    }
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();
}