Ascend C 学習ノート

AI プロセッサ#

構成#

HBM（ストレージ）
DVPP（ハードウェアのエンコードとデコード）
CPU（4 コア、ARM）：命令の配布などの機能
AI Core NPU（AI 計算）：AI 計算、行列計算（16x16x16）、スカラー計算、ベクトル計算（128）などを実行します。

一般的な計算と AI 計算の違い#

一般的な計算：一般的な計算タスクに適用されます。加算、減算、乗算、除算などです。
例えば、行列の乗算を計算する場合、直列計算しかできず、速度が遅いです。

NPU 計算では、ベクトル計算ユニットを使用して複数のデータを同時に計算することができ、速度が速くなります。または、行列計算ユニットを直接使用して行列の乗算を計算することもできます。

ヘテロジニアス計算：CPU + NPU、CPU は制御を担当し、NPU は計算を担当します。

Ascend C#

ネイティブオペレータの開発、自動スケジューリング。

グラフの最適化、自動パイプライン化。
オペレータのディープフュージョン、異なるオペレータを自動的に統合し、データメモリのコピーを減らす。

計算グラフの実行の沈降、すべてのオペレータを NPU 上で実行する。

自動チューニング AOE

Ascend C#

AI Core#

計算ユニット、ストレージユニット、制御ユニット。

20240711102909

内部構造：
スカラー計算ユニット（CPU）が命令を送信します。
20240711103146

テンソルの概念
名前
形状
データタイプ
フォーマット

フォーマット#

N：バッチサイズ、画像の数
H：高さ
W：幅
C：チャンネル数、RGB は 3

Ascend C の利点#

20240711105258

デバイスモジュール#

ホストとデバイス

ホストはサーバーを指し、デバイスは NPU を指します。

Ascend C オペレータの開発#

Add オペレータの開発#

オペレータの登録#

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();
}

ここで、私たちのカーネルオペレータ関数 add_custom を定義しました。この関数は、x、y、z の 3 つのパラメータを受け取ります。これらのパラメータは、入力テンソル 2 つと出力テンソルです。

GM_ADDR は__gm__ uint8_t*の別名です。グローバルメモリ上のアドレスを表します。

オペレータの実装#

ここでは、オペレータを実装するためにクラスを使用しました。

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

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        ...
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        ...
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        ...
    }

private:
    ...
};

クラスでは、Init と Process の 2 つの主要な関数を公開しています。

__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));
}

Init 関数は、メモリアドレスを割り当てて初期化します。

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

Process 関数は、CopyIn、Compute、CopyOut の 3 つのパイプライン構成の関数をループで呼び出します。

CopyIn 関数は、データをグローバルメモリからローカルメモリにコピーします。Compute 関数は計算を行い、CopyOut 関数は計算結果をローカルメモリからグローバルメモリにコピーします。

__aicore__ inline void CopyIn(int32_t progress)
{
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
    DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
    DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}
__aicore__ inline void Compute(int32_t progress)
{
    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<half>(zLocal);
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}
__aicore__ inline void CopyOut(int32_t progress)
{
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    outQueueZ.FreeTensor(zLocal);
}

計算では、ベクトル加算を使用してこのオペレータを実装しました。

すべての命令は並行して実行されるため、このオペレータではキューを使用してデータ依存性を確保しています。

DeQue は、必要なデータがキューにあるまでブロックされます。

オペレータの呼び出し#

カーネルオペレータを呼び出す際には、

<<<...>>> を使用して関数を呼び出します。

blockDim：デバイスによって異なるブロックサイズを設定します。
l2ctrl：現在は nullptr
stream：全体のストリームを表します。

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);
}

Sinh オペレータの開発#

Sinh の他の構造は Add と似ていますが、主な違いはパラメータの数です。

主な違いは Compute 関数の実装です。

ここでは、Muls Exp を使用して、まず $e^{-x}$ を求めます。
次に $e^x$ を求めます。
最後にそれらを引いて 2 で割ることで sinh の結果が得られます。

__aicore__ inline void Compute(int32_t progress)
{
    LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
    LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();
    LocalTensor<DTYPE_Y> tmpLocal = tmpBuffer.AllocTensor<DTYPE_Y>();
    // e^-x
    Muls(tmpLocal, xLocal, static_cast<DTYPE_X>(-1), this->tileLength);
    Exp(tmpLocal, tmpLocal, this->tileLength);
    // e^x 
    Exp(yLocal, xLocal, this->tileLength);
    // (e^x - e^-x)
    Sub(yLocal, yLocal, tmpLocal, this->tileLength);
    // (e^x - e^-x) / 2
    Muls(yLocal, yLocal, static_cast<DTYPE_X>(0.5), this->tileLength);
    outQueueY.EnQue<DTYPE_Y>(yLocal);
    inQueueX.FreeTensor(xLocal);
    tmpBuffer.FreeTensor(tmpLocal);
}