Kirin 9020 series processors
Kirin 9030 series processors
Kirin X90 series processors
Do not create or initialize TPipe inside an object. If TPipe is created inside an object, the optimization of constants in the object by the compiler may be affected, causing scalar performance deterioration.
Only one TPipe object can exist globally at a time. If multiple TPipe objects are defined at the same time, random behaviors such as suspension may occur. If multiple TPipe objects are required, call Destroy to release the previous TPipe object first.
template <typename ComputeT> class KernelExample {
public:
__aicore__ inline KernelExample() {}
__aicore__ inline void Init(..., TPipe* pipeIn)
{
// ...
pipe = pipeIn;
pipe->InitBuffer(xxxBuf, BUFFER_NUM, xxxSize);
// ...
}
private:
// ...
TPipe* pipe;
// ...
};
extern "C" __global__ __aicore__ void example_kernel(...) {
// ...
TPipe pipe;
KernelExample<float> op;
op.Init(..., &pipe);
// ...
}