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 或跳过零权重,需要新的明确调度,同时保持数值不变。
可运行文件
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_acc、out_code 和 out_activation,同时比较 out_row/out_col 及已接收的输出传送。第四个结果被接收后,done 上升;此时 pooled 有效。接收十六个输入不意味着完成。此 RTL 使用九个并行乘积和两个寄存级,不同于之前的双 PE、24 周期估计。频率/功耗需要综合或器件证据。
运行共享的自检查测试平台:
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%。相同频率、供数能力和合并成本仍需要独立实现证据。