Capstone: a small inference path and report

Goal and acceptance criteria

Keep model.mjs from the preceding lab,create capstone-lab.mjs beside it and run node capstone-lab.mjs. Connect one convolution layer,bias,requantization,ReLU and max pooling. Report values separately from cost. This is not a complete classifier or training engine. You may manually inspect the files through a supported installed-Digital-Design-Studio file workflow; execution inside DDS is not established.

Finish when direct convolution and patch-dot accumulators agree,intermediate/output fixtures match,and your report includes shape,MACs,assumed cycles,slot utilization and unique tensor bytes. Add one independently hand-calculated changed-input fixture. Keep the command,file version and assumptions with the log. Investigate a failed comparison instead of merely editing its expected value.

Worked reference

Input is 4×4 with values 1 through 16. Kernel [1,0,−1;1,0,−1;1,0,−1] subtracts the right sample from the left in each row. Each row gives −2, so every window sums to −6. Bias eight gives four accumulators equal to two. Multiplier 0.5 and zero point zero produce code one. ReLU retains it and 2×2 pooling gives [1]. Retain asymmetric boundary fixtures too; equal outputs can conceal ordering mistakes.

Shape [M,Ho,Wo] is [1,2,2]. Count 36 MACs, including zero-weight terms. With two single-MAC PEs sharing each output reduction, assume five MAC cycles plus one combine/write cycle per output, no overlap or stalls. Modeled total is 24 cycles. MAC-only slot utilization is 36/(4×5×2)=90%. This is not measured device latency.

Unique int8 tensors use input 16, weights 9 and convolution outputs 4 bytes, plus a four-byte bias. Pooled output uses one byte. Actual traffic may include reloads and partial sums. Scales sx=sw=1, sy=2 explain multiplier 0.5.

Changing bias to zero predicts accumulator −6, code −3, activated zero and pooled zero. Changing only the first input from one to two gives first accumulator three, tie 1.5 rounding to two, three remaining outputs one and pooled two. Test changes separately. More PEs or skipping zero weights requires a new explicit schedule while preserving values.

Runnable file

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 capstone: integer convolution

Attachments cnn_window3x3.sv and cnn_conv3x3.sv connect windows to nine signed 8×8 products, a signed 24-bit sum/bias, nearest-even division by two, int8 saturation, ReLU and a 2×2 maximum. Geometry is the preceding 4×4, single-channel, valid-3×3 contract. Input/weight scales are one, output scale two and all zero points zero. Default kernel is [1,0,−1;1,0,−1;1,0,−1], bias eight. Kernel parameters may change; recheck the 24-bit sum bound.

Compare out_acc, out_code and out_activation separately, with out_row/out_col and accepted output transfers. done rises after the fourth result is accepted; pooled is then valid. Sixteen accepted inputs do not imply completion. This RTL uses nine parallel products and two registered stages, rather than the earlier two-PE,24-cycle estimate. Frequency/power require synthesis or device evidence.

Run the shared self-checking bench:

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

Default accumulators are two, codes/activations one and pooled output one. Checks include signed boundaries, saturation, positive/negative ties, a dense signed kernel, pauses and reset. Verify PASS frames=6 windows=24 outputs=24 kernels=2 in your actual log. Change the first input to two: first sum three rounds to code two and pooled becomes two. Then choose another signed kernel, hand-calculate its first window, update the independent coefficient function and rerun every fixture. Do not copy DUT results into the golden model.