The core of Ascend C’s vector model lies in decomposing computation into three tightly coupled, asynchronous stages: CopyIn, Compute, and CopyOut. These stages operate across distinct memory domains—Global Memory (off-chip) and Local Memory (on-chip AI Core)—and are coordinated via abstracted queuing primitives. The paradigm abstracts hardware topology through TPosition and TQue, enabling developers to focus on data flow rather than register allocation or cache hierarchy.

Operator Design Workflow

Developing a custom vector operator begins with formal specification:

• Mathematical Definition: For an addition operator, define z = x + y, where all operands are tensors of type half (float16).
• Data Shapes: Input tensors x and y have shape (8, 2048), output z matches this shape in ND format.
• API Selection: Use DataCopy for memory transfers, Add for vectorized arithmetic, and TQue for inter-stage synchronization.

Kernel Structure and Execution Model

Each kernel is a single AI Core instance, launched via __global__ __aicore__ and invoked with <<<blockDim, l2ctrl, stream>>>. The kernel function instantiates an operator class that encapsulates state and logic:

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

#ifndef __CCE_KT_TEST__
void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z) {
    add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}
#endif

The Init() method configures memory mappings and buffer resources. Process() orchestrates the pipeline by repeatedly invoking the three stages.

Data Flow and Memory Abstraction

Ascend C uses GlobalTensor to reference off-chip data and LocalTensor for on-chip scratch space. Communication between pipeline stages occurs through TQue queues, bound to logical positions: TPosition::VECIN and TPosition::VECOUT.

The data path operates as follows:

• CopyIn: Copy data from GlobalTensor to LocalTensor using DataCopy, then enqueue into inQueueX and inQueueY.
• Compute: Dequeue inputs, perform vectorized addition via Add(), enqueue result to outQueueZ, and free input buffers.
• CopyOut: Dequeue result from outQueueZ, copy to output GlobalTensor, and release local buffer.

Memory management is handled by TPipe, which allocates buffer space for queues using InitBuffer(). The buffer size is specified in bytes, not element count, allowing fine-grained control over memory footprint.

Operator Class Implementation

The operator class defines public initialization and processing interfaces, and private pipeline methods:

class KernelAdd {
public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z);
    __aicore__ inline void Process();

private:
    __aicore__ inline void CopyIn(int32_t step);
    __aicore__ inline void Compute(int32_t step);
    __aicore__ inline void CopyOut(int32_t step);

private:
    TPipe pipe;
    TQue<TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ;
    GlobalTensor<half> xGm, yGm, zGm;
};

Init() Implementation:

__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

    pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
}

Each AI Core processes a unique segment of the global tensor, determined by its block index. Buffer sizes are configured based on tile and buffer dimensions.

Process() and Pipeline Stages:

__aicore__ inline void Process() {
    constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; ++i) {
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

__aicore__ inline void CopyIn(int32_t step) {
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();

    DataCopy(xLocal, xGm[step * TILE_LENGTH], TILE_LENGTH);
    DataCopy(yLocal, yGm[step * TILE_LENGTH], TILE_LENGTH);

    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}

__aicore__ inline void Compute(int32_t step) {
    LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();

    Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    outQueueZ.EnQue(zLocal);

    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

__aicore__ inline void CopyOut(int32_t step) {
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    DataCopy(zGm[step * TILE_LENGTH], zLocal, TILE_LENGTH);
    outQueueZ.FreeTensor(zLocal);
}
</half></half></half></half></half></half></half></half></half></half></half></half>

Each stage operates on a single tile of data. The loop count is derived from the product of tile count and buffer depth, ensuring full utilization of double-buffering.

Data Partitioning Strategy

Ascend C employs SPMD (Single Program, Multiple Data) execution. The total dataset is divided across logical cores, not physical ones. Each core handles a BLOCK_LENGTH segment. Within each core, data is further partitioned into TILE_NUM tiles, each processed independently.

To maximize throughput, double buffering is enabled via BUFFER_NUM = 2, allowing one buffer to be filled while another is being processed. This overlaps data transfer with computation, hiding latency.

constexpr int32_t TOTAL_LENGTH = 8 * 2048;
constexpr int32_t USE_CORE_NUM = 8;
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;
constexpr int32_t TILE_NUM = 8;
constexpr int32_t BUFFER_NUM = 2;
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // 128 elements

With this configuration, each tile holds 128 half-precision values. The 8-tile, 2-buffer layout ensures that at any moment, two tiles are in flight: one being copied in, one being computed, and one being copied out—creating a continuous pipeline.

Performance hinges on aligning tile size with vector instruction width and minimizing buffer contention. Larger tiles reduce enqueue/dequeue overhead but increase memory pressure. Smaller tiles improve pipelining granularity but increase synchronization cost.