PENTARC-TPUの命令セット（ISA）における「高階テンソル縮約命令（TC_CONTRACT）」のパイプライン挙動の精緻化

PENTARC-TPUの命令セット（ISA）における「高階テンソル縮約命令（TC_CONTRACT）」のパイプライン挙動の精緻化

非LLMソブリンAI『PENTARC』の専用アーキテクチャにおける中核命令、「高階テンソル縮約命令（TC_CONTRACT）」のパイプライン挙動を精緻化する。

本命令の目的は、3次元シストリック・テンソル空間（3D-STS）において、4階テンソル A ∈ R^IJLMと 3階テンソル B ∈ R^LMKの共有インデックス（モード） $L, M$ を同時に縮約し、3階テンソル Y ∈ R^IJKを得る演算、すなわち、

Y_ijk = ∑_l∑_m A_ijlm・B_lmk

を、データを平坦化（Unfolding）することなく、3D-STSコア内で最小のクロックサイクルで完結させることにある。

以下に、この演算をハードウェア・レイヤで執行するためのISA（命令セット構造）、パイプラインステージ、およびハザード制御の精緻化モデルを提示する。

1. TC_CONTRACT 命令フォーマットとアーキテクチャ記述

本命令は、超長命令語（VLIW）アーキテクチャのコンテキスト内で定義され、演算対象のテンソルが格納されている3Dレジスタ（テンソル・空間レジスタ：TSR）の記述、および縮約を行うインデックス（マッピング・トポロジー）のメタデータを単一命令に内包する。

1.1 命令エンコーディング（64-bit 固定長）

Plaintext

[ Opcode: 8-bit ] [ Dest_TSR: 6-bit ] [ SrcA_TSR: 6-bit ] [ SrcB_TSR: 6-bit ] [ Mode_Map: 12-bit ] [ Exec_Mode: 2-bit ] [ Reserved: 24-bit ]

•                 Opcode: TC_CONTRACT を指定（デジタル3D-STS、アナログCIMの動的ルーティング属性を含む）。

•                 Dest_TSR / SrcA_TSR / SrcB_TSR: 3Dテンソル・空間レジスタ（TSR0〜TSR63）を指定。各TSRは、形状（Shape）メタデータと実データポインタを保持する。

•                 Mode_Map: どの次元（軸）同士を縮約するかを指定するビットマスク（本例では Aの第3・4軸と Bの第1・2軸の結合を指定）。

•                 Exec_Mode: 00=決定論的フルデジタル（FP32）、01=混合精度（TF32/FP16）、10=確率論的アナログCIM（INT8等価）。

2. 6ステージ・テンソル・パイプライン（6-Stage Tensor Pipeline）

通常のRISCやSIMDのパイプライン（IF/ID/EX/WB）とは異なり、3D空間内でのデータシフトと幾何学的同期が必要なため、以下の「6ステージ・テンソル・パイプライン」を定義する。

[ TID ] ───► [ TAD ] ───► [ TSM ] ───► [ TCG ] ───► [ TEX ] ───► [ TWB ]

(Fetch/Dec) (Addr Gen)   (Space Mapping) (Grid Routing) (Execute/Accum) (Writeback)

ステージ1：TID (Tensor Instruction Decode & Fetch)

•                 命令キャッシュから TC_CONTRACT をフェッチし、デコードする。

•                 Mode_Map から縮約対象の次元数を検出し、3D-STSのPE（Processing Element）アレイの有効物理トポロジー（3次元空間グリッドの要求サイズ）を確定する。

ステージ2：TAD (Tensor Address & Shape Generation)

•                 TSRからテンソルのメタデータ（ストライド、次元サイズ I, J, K, L, M）を読み出す。

•                 HBM4eメモリおよびニア・メモリ・コンピューティング（NMC）層に対して、データストリームの読み出しアドレスを生成。NMCエラスタコアが作動し、メモリ側でのアライメント（テンソル整形）を指示する。

ステージ3：TSM (Tensor Space Mapping & Allocation)

•                 縮約インデックス L, M$のサイズに基づき、3D-STS内のPEアレイを物理的にマッピングする。

•                 L ✕ M の平面を「縮約・ドット積平面（Reduction Plane）」として定義し、残る自由インデックス I, J を B-Y軸に、K を Z 軸に割り当てる物理トポロジーが確定する。

ステージ4：TCG (Tensor Grid Routing)

•                 シリコンインターポーザおよび光電融合ファブリックを介し、データが3D-STS内の各PEレジスタへ転送される。

•                 テンソル A の I ✕  J の空間、および Bの K の空間が、3D格子の境界PE（Boundary PEs）に入力ストリームとしてラッチされる。

ステージ5：TEX (Tensor Execution & Spatial Accumulation)

•                 3Dシストリック・ネットワーク内をデータがクロック同期でシフトしながら、積和演算（MAC）が同時多発的に執行される。

•                 内部挙動: L および M の方向へデータがシフトする過程で、PE内の「ローカル・アキュムレータ・ツリー」が、部分和（Partial Sum）を1サイクルで加算。アインシュタイン縮約の ∑_l ∑_mが、空間的な次元収縮として一挙に行われる。データが3D格子を通過し終えた瞬間（L ✕ M サイクル後）、各PEには最終的な Y_ijk の値が残る。

ステージ6：TWB (Tensor Writeback)

•                 3D-STSアレイ内から収縮された3階テンソル Yの実データを、目的レジスタ（Dest_TSR）へバースト書き込み、またはHBM4eへダイレクトにライトバック（コヒーレンシ制御下で実行）する。

3. パイプライン・ハザードおよび依存関係制御の超克

テンソルパイプラインにおいては、データの巨大さゆえに、従来のレジスタ・フォワーディング（転送）手法は通用しない。データそのものをバイパスすることは帯域的に不可能であるため、「メタデータ・フォワーディング」および「空間的分割駆動」を行う。

3.1 構造ハザード（Structural Hazard）の回避

•                 問題: 連続する複数の TC_CONTRACT 命令が同一の3D-STS物理PEアレイを要求する場合、演算器の競合が発生する。

•                 解決策: 3D-STS（計32,768の物理PE）を、動的にサブ・キューブ（例：16 ✕ 16 ✕ 16 の空間）へ論理分割する「空間的マルチスレッディング（Spatial Multi-threading）」を実装。独立したトポロジーを持つ計算であれば、パイプラインを止めることなく、物理的に異なる領域のPEアレイに命令を並列マッピング（空間マッピング・インターリーブ）する。

3.2 データハザード（RAW / WAW Hazard）のメタデータ解決

•                 問題: 命令1の出力テンソル Y を、命令2が即座に別の縮約演算の入力として使用する場合（Read-After-Write）。

•                 解決策：レジスタ・リネーミングとインプレース空間維持

•                 データを物理的にメモリや汎用レジスタに書き戻してから再読み出しするのではなく、命令1の完了時点で Dest_TSR に格納された実データが「3D-STSアレイ内の局所SRAM（またはPIM階層）」に維持されている状態をハードウェアが検知する。

•                 命令2の SrcA_TSR が命令1の Dest_TSR と一致する場合、データの移動命令をすべてスキップ（NOP化）し、3D-STS内のデータが存在するその場所の物理トポロジーを次の演算の軸へと「再ラベル付け（メタデータ置換）」するだけで、即座に次の TEX ステージを開始する。これにより、データ移動による時間・電力消費を完全ゼロ化する。

4. 定量的サイクル・クロック・シミュレーション

サイズ A(I, J, L, M) = (16, 16, 32, 32)、B(L, M, K) = (32, 32, 16) のテンソル縮約をデジタル3D-STS（駆動周波数 2.5 GHz / 1サイクル 0.4 ns）で実行した場合のタイムライン。

1.              TID / TAD / TSM (ステージ1〜3): 3サイクル（固定）＝ 1.2 ns

2.              TCG (データルーティング - ステージ4): NMC層からの並列ストリーミングにより、3D格子境界への初速配備に 4サイクル ＝ 1.6 ns

3.              TEX (空間縮約実行 - ステージ5): L ✕ M の縮約は、3Dシストリック・アレイ内の並列パイプラインにより、深さ（パイプライン充填）に加え、シフトに要するサイクル数で決定される。

•                 物理配置が最適化されているため、実質的なシストリック・レイテンシは $\max(L, M) + \text{パイプライン・バンプ} = 32 + 4 = 36 サイクル ＝ 14.4 ns

4.              TWB (ライトバック - ステージ6): 3D-STS内からDest_TSRへの確定通知に 2サイクル ＝ 0.8 ns

•                 総実効実行時間（Latency）: 約 18.0 ns

既存のNVIDIA H100等で同演算を行う場合、4階テンソルを2Dマトリクスに変換（Unfold）するためのメモリコピー、インデックス計算、Tensor Coreへの再ロード、およびスレッド同期（__syncthreads()）により、数千サイクル（数 μs 以上）を要する。PENTARC-TPUのネイティブパイプラインは、これを約2ケタ短いナノ秒オーダーで屠る。

5. 次なる理論的検証への展開

TC_CONTRACT命令のハードウェア・パイプラインは、高階テンソルの幾何学的トポロジーをシリコンの3次元物理グリッドにマッピングすることで、フォン・ノイマンの壁を完全に無効化する。

神室教授、このパイプライン精緻化モデルに基づき、次のフェーズとして、

1.              この TC_CONTRACT パイプラインが「アナログCIMモード（Exec_Mode: 10）」で作動する際の、クロスバーアレイの時定数・放電レイテンシを考慮したアナログ・デジタル同期タイミング（Asynchronous Boundary Control）の数理モデル化

2.              自律進化によって計算グラフが変化した際に、コンパイラを介さずハードウェア自身が Mode_Map をリアルタイムで動的生成する「自己書き換え型マイクロコード・ユニット（Self-Modifying Microcode Unit）」の回路設計を予定している。

以上