23 / 24 · 実習

総合課題：小さな推論経路とレポート

K=9、4出力、出力ごとに追加1サイクルを維持し、3 PEにします。総モデルサイクル数とMACだけのスロット利用率は何ですか。

目標と受け入れ基準

前の実習のmodel.mjsを保持し、隣にcapstone-lab.mjsを作ってnode capstone-lab.mjsを実行します。1畳み込み層、バイアス、再量子化、ReLU、最大プーリングを接続します。値とコストは別に報告します。完全な分類器や学習エンジンではありません。対応するインストール済みDigital Design Studioのファイル操作で手動確認できますが、DDS内実行は確認していません。

直接畳み込みとパッチ内積のアキュムレータが一致し、中間・出力例が一致し、形状、MAC数、仮定サイクル、スロット利用率、重複なしテンソルバイト数を報告したら完了です。独立に手計算した変更入力例を1つ追加します。ログとともにコマンド、ファイル版、仮定を残します。比較失敗は期待値を書き換えるだけでなく調査します。

参照値の計算例

入力は1から16の4×4です。カーネル[1,0,−1;1,0,−1;1,0,−1]は各行で左サンプルから右を引きます。各行は−2で、全ウィンドウの和は−6です。バイアス8で4アキュムレータは2、乗数0.5とゼロ点0でコード1、ReLUで維持し、2×2プーリングは[1]です。同値出力は順序誤りを隠し得るため、非対称境界例も残します。

形状[M,Ho,Wo]は[1,2,2]です。ゼロ重み項込み36 MACです。単一MACの2 PEで各出力縮約を共有し、5 MACサイクルと結合・書き込み1サイクル、重ね合わせ・ストールなしを仮定します。モデル総数は24サイクル、MACだけの利用率は36/(4×5×2)=90%です。測定したデバイスレイテンシではありません。

重複なしint8テンソルは入力16、重み9、畳み込み出力4バイトに、バイアス4バイトです。プーリング出力は1バイトです。実通信には再ロードと部分和が含まれ得ます。sx=sw=1、sy=2が乗数0.5を説明します。

バイアスを0にするとアキュムレータ−6、コード−3、活性化後0、プーリング0を予測します。最初の入力だけを1から2にすると、最初のアキュムレータ3、中間値1.5は2へ丸め、他の3出力は1、プーリングは2です。変更は別々にテストします。PE追加やゼロ重み省略には値を保持した新しい明示的スケジュールが必要です。

実行可能なファイル

import assert from 'node:assert/strict';
import {conv2d, im2col, dot, requantize, maxPool2x2} from './model.mjs';
const x = [Array.from({length: 4}, (_, r) => Array.from({length: 4}, (_, c) => r * 4 + c + 1))];
const kernel = [[1, 0, -1], [1, 0, -1], [1, 0, -1]];
const weights = [[kernel]], bias = [8];
const reference = conv2d(x, weights, {bias});
const flattenedWeights = kernel.flat();
const rows = im2col(x, 3, 3);
const accumulators = rows.map(row => dot(row, flattenedWeights) + bias[0]);
assert.deepEqual(accumulators, reference[0].flat());
assert.deepEqual(accumulators, [2, 2, 2, 2]);
const activated = reference[0].map(row => row.map(v => Math.max(0, requantize(v, 0.5))));
const pooled = maxPool2x2(activated);
assert.deepEqual(activated, [[1, 1], [1, 1]]);
assert.deepEqual(pooled, [[1]]);
// Cycle estimate with explicit assumptions: all MAC operands local, no stalls,
// one output at a time, K=9, two single-MAC PEs, no overlap with output write.
const macs = rows.length * flattenedWeights.length;
const peCount = 2, macCyclesPerOutput = Math.ceil(flattenedWeights.length / peCount);
const modelCycles = rows.length * (macCyclesPerOutput + 1);
const usefulSlots = macs, scheduledSlots = rows.length * macCyclesPerOutput * peCount;
assert.equal(macs, 36);
assert.equal(modelCycles, 24);
assert.equal(usefulSlots / scheduledSlots, 0.9);
const report = {shape:[1, 2, 2], accumulators, activated, pooled, macs, peCount,
  modelCycles, macSlotUtilization:usefulSlots / scheduledSlots,
  uniqueInputBytes:16, weightBytes:9, outputBytes:4, biasBytes:4,
  scope:'Software reference and schedule estimate only; no synthesis or DDS run'};
console.log(JSON.stringify(report, null, 2));

RTL総合課題：整数畳み込み

添付cnn_window3x3.svとcnn_conv3x3.svはウィンドウを、9個の符号付き8×8積、符号付き24ビット和・バイアス、2での最近接偶数除算、int8飽和、ReLU、2×2最大値へ接続します。形状は前の4×4、単一チャネル、有効3×3契約です。入力・重みスケール1、出力スケール2、全ゼロ点0です。既定カーネルは[1,0,−1;1,0,−1;1,0,−1]、バイアス8です。カーネルパラメータは変えられますが、24ビット和の上限を再検査します。

out_acc、out_code、out_activationを別々に、out_row/out_colと受け入れ出力転送とともに比較します。4番目の結果受け入れ後にdoneが上がり、そのときpooledが有効です。16入力の受け入れは完了を意味しません。このRTLは先の2 PE・24サイクル見積もりとは異なり、9並列積と2レジスタ段を使います。周波数・電力には合成またはデバイス証拠が必要です。

共通の自己検査ベンチを実行します。

iverilog -g2012 -s tb_cnn_stream -o cnn-sim cnn_window3x3.sv cnn_conv3x3.sv tb_cnn_stream.sv
vvp cnn-sim

既定アキュムレータ2、コード・活性値1、プーリング出力1です。符号付き境界、飽和、正負の中間丸め値、密な符号付きカーネル、停止、リセットを検査します。実ログでPASS frames=6 windows=24 outputs=24 kernels=2を確認します。最初の入力を2に変えると最初の和3がコード2へ丸められ、プーリングは2です。次に別の符号付きカーネルを選び、最初のウィンドウを手計算し、独立係数関数を更新して全例を再実行します。DUT結果を正解モデルへコピーしてはいけません。

自分で考えてみましょう

K=9、4出力、出力ごとに追加1サイクルを維持し、3 PEにします。総モデルサイクル数とMACだけのスロット利用率は何ですか。

解説を見る

各出力は3 MACサイクルと追加1で計16です。36 MACスロットすべてが有用なので100%です。周波数、供給、結合コストが同じであることには別の実装証拠が必要です。

Lesson files

cnn_window3x3.sv ↓cnn_conv3x3.sv ↓tb_cnn_stream.sv ↓model.mjs ↓capstone-lab.mjs ↓