Итоговый проект: малый путь инференса и отчёт

Цель и критерии приёмки

Сохраните model.mjs из предыдущей работы,создайте рядом capstone-lab.mjs и выполните node capstone-lab.mjs. Соедините один слой свёртки,bias,повторное квантование,ReLU и максимальный пулинг. Сообщайте значения отдельно от затрат. Это не полный классификатор или движок обучения. Можно вручную исследовать файлы через поддерживаемую работу с файлами установленной Digital Design Studio; выполнение внутри DDS не установлено.

Завершите, когда аккумуляторы прямой свёртки и оконных скалярных произведений согласуются,промежуточные/выходные примеры совпадают,а отчёт включает форму,MAC,предполагаемые такты,загрузку мест и уникальные байты тензоров. Добавьте один независимо рассчитанный вручную пример изменённого входа. Сохраните команду,версию файла и допущения с журналом. Исследуйте ошибочное сравнение вместо простого изменения ожидаемого значения.

Разобранный эталон

Вход — 4×4 со значениями 1–16. Ядро [1,0,−1;1,0,−1;1,0,−1] вычитает правый отсчёт из левого в каждой строке. Каждая строка даёт −2, поэтому каждое окно даёт сумму −6. Bias восемь даёт четыре аккумулятора, равных двум. Множитель 0.5 и нулевая точка ноль дают код один. ReLU сохраняет его, пулинг 2×2 даёт [1]. Сохраняйте также асимметричные граничные примеры; равные выходы могут скрыть ошибки порядка.

Форма [M,Ho,Wo] — [1,2,2]. Считайте 36 MAC, включая члены нулевых весов. Для двух PE по одному MAC, разделяющих свёртку каждого выхода, предположите пять тактов MAC плюс один объединения/записи на выход без перекрытия или остановок. Моделируемый итог — 24 такта. Загрузка мест только MAC — 36/(4×5×2)=90%. Это не измеренная задержка устройства.

Уникальные тензоры int8 используют вход 16, веса 9 и выходы свёртки 4 байта плюс четырёхбайтный bias. Выход пулинга использует байт. Реальный трафик может включать перезагрузки и частичные суммы. Масштабы sx=sw=1, sy=2 объясняют множитель 0.5.

Изменение bias на ноль предсказывает аккумулятор −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-битной суммой/bias, делением на два с округлением к ближайшему чётному, насыщением int8, ReLU и максимумом 2×2. Геометрия соответствует предыдущей спецификации 4×4, одного канала и 3×3 без дополнения. Масштабы входа/весов — один, выходной — два, все нулевые точки — ноль. Ядро по умолчанию — [1,0,−1;1,0,−1;1,0,−1], bias восемь. Параметры ядра можно менять; перепроверяйте границу 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 в реальном журнале. Измените первый вход на два: первая сумма три округляется до кода два, pooled становится два. Затем выберите другое знаковое ядро, рассчитайте первое окно вручную, обновите независимую функцию коэффициентов и повторите каждый пример. Не копируйте результаты DUT в золотую модель.