Altifigence Academy

23 / 24 · 實作

總整實作:小型推論路徑與報告

保持 K=9、四個輸出與各輸出額外一週期,改用三個 PE。模型總週期與僅計 MAC 的槽利用率各是多少?

目標與驗收條件

保留前一實作的 model.mjs,在旁邊建立 capstone-lab.mjs,執行 node capstone-lab.mjs。連接一個卷積層、偏置、再量化、ReLU 與最大池化。數值與成本分別回報。這不是完整分類器或訓練引擎。可透過已安裝 Digital Design Studio 支援的檔案流程手動檢視檔案;未證實在 DDS 內執行。

當直接卷積與視窗內積累加器一致、中間與輸出測試資料符合,且報告包含形狀、MAC、假設週期、槽利用率與唯一張量位元組數時,才完成。加入一組獨立手算的變更輸入測試資料。紀錄旁保留命令、檔案版本與假設。比較失敗時先調查,不要只修改預期值。

參考計算範例

輸入為包含 1 到 16 的 4×4。核心 [1,0,−1;1,0,−1;1,0,−1] 在各列以左樣本減右樣本。各列得到 −2,因此每個視窗總和為 −6。偏置八,使四個累加器都等於二。乘數 0.5、零點零產生編碼一。ReLU 保留它,2×2 池化得到 [1]。也保留非對稱邊界測試資料;相同輸出可能掩蓋順序錯誤。

形狀 [M,Ho,Wo] 為 [1,2,2]。計算 36 次 MAC,包括零權重項。兩個單 MAC PE 共用各輸出歸約,假設每個輸出五個 MAC 週期加一個合併與寫入週期,沒有重疊或停頓。模型總計 24 週期。僅計 MAC 的槽利用率為 36/(4×5×2)=90%。這不是裝置延遲量測。

唯一 int8 張量使用:輸入 16、權重 9、卷積輸出 4 位元組,加上一個四位元組偏置。池化輸出使用一位元組。實際流量可能包括重新載入與部分和。尺度 sx=sw=1、sy=2 解釋乘數 0.5。

偏置改為零,預測累加器 −6、編碼 −3、活化零與池化零。只將第一個輸入從一改為二,第一累加器為三,同距 1.5 捨入為二,其餘三個輸出為一,池化為二。分別測試變更。更多 PE 或跳過零權重,需要新的明確排程,同時保留數值。

可執行檔案

javascript
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 將視窗連至九個有號 8×8 乘積、有號 24 位元總和與偏置、最近偶數除以二、int8 飽和、ReLU 與 2×2 最大值。幾何沿用前述 4×4、單通道、無填補有效 3×3 契約。輸入與權重尺度為一,輸出尺度二,所有零點為零。預設核心為 [1,0,−1;1,0,−1;1,0,−1],偏置八。核心參數可以改變;重新檢查 24 位元總和上界。

搭配 out_row/out_col 與被接受的輸出傳輸,分別比較 out_acc、out_code 與 out_activation。第四個結果被接受後,done 上升;此時 pooled 有效。十六個被接受輸入不代表完成。此 RTL 使用九個平行乘積與兩個暫存器級,不同於前述兩個 PE、24 週期估計。頻率與功耗需要合成或裝置證據。

執行共用的自我檢查測試平台:

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

預設累加器為二,編碼與活化值為一,池化輸出為一。檢查包括有號邊界、飽和、正負同距、密集有號核心、暫停與重設。在實際紀錄中驗證 PASS frames=6 windows=24 outputs=24 kernels=2。第一個輸入改為二:第一總和三捨入為編碼二,池化變為二。再選另一個有號核心,手算其第一視窗,更新獨立係數函式,並重跑所有測試資料。不要將 DUT 結果複製至黃金模型。

自己試試看

保持 K=9、四個輸出與各輸出額外一週期,改用三個 PE。模型總週期與僅計 MAC 的槽利用率各是多少?

查看解說

各輸出需要三個 MAC 週期加一個額外週期,因此總計十六。全部 36 個 MAC 槽都有用,利用率為 100%。相同頻率、供應與合併成本仍需要獨立實作證據。

Lesson files

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

此選擇適用於本瀏覽器,隨時可從頁尾變更。