Progetto finale: piccolo percorso di inferenza e rapporto

Obiettivo e criteri di accettazione

Mantieni model.mjs dal laboratorio precedente,crea capstone-lab.mjs accanto ed esegui node capstone-lab.mjs. Collega uno strato di convoluzione,bias,riquantizzazione,ReLU e max pooling. Riporta i valori separatamente dal costo. Non è un classificatore completo o un motore di addestramento. Puoi ispezionare manualmente i file tramite un flusso file supportato di Digital Design Studio installato; l'esecuzione dentro DDS non è attestata.

Termina quando gli accumulatori della convoluzione diretta e dei prodotti scalari delle patch concordano,i casi intermedi/di uscita corrispondono,e il rapporto include forma,MAC,cicli ipotizzati,utilizzo delle posizioni e byte unici dei tensori. Aggiungi un caso con ingresso modificato calcolato indipendentemente a mano. Conserva comando,versione del file e ipotesi insieme al log. Indaga un confronto fallito anziché limitarti a modificarne il valore atteso.

Riferimento svolto

L'ingresso è 4×4 con valori da 1 a 16. Il kernel [1,0,−1;1,0,−1;1,0,−1] sottrae il campione destro da quello sinistro in ogni riga. Ogni riga dà −2, quindi ogni finestra somma a −6. Il bias otto dà quattro accumulatori uguali a due. Moltiplicatore 0.5 e punto zero zero producono codice uno. ReLU lo mantiene e il pooling 2×2 dà [1]. Conserva anche casi asimmetrici ai confini; uscite uguali possono nascondere errori di ordinamento.

La forma [M,Ho,Wo] è [1,2,2]. Conta 36 MAC, inclusi i termini con pesi zero. Con due PE da un MAC che condividono ogni riduzione di uscita, presumi cinque cicli MAC più un ciclo di combinazione/scrittura per uscita, senza sovrapposizione o stalli. Il totale modellato è 24 cicli. L'utilizzo delle posizioni dei soli MAC è 36/(4×5×2)=90%. Non è latenza misurata sul dispositivo.

I tensori int8 unici usano ingresso 16, pesi 9 e uscite di convoluzione 4 byte, più un bias di quattro byte. L'uscita del pooling usa un byte. Il traffico reale può includere ricaricamenti e somme parziali. Le scale sx=sw=1, sy=2 spiegano il moltiplicatore 0.5.

Cambiare il bias a zero prevede accumulatore −6, codice −3, valore attivato zero e pooling zero. Cambiare soltanto il primo ingresso da uno a due dà primo accumulatore tre, equidistanza 1.5 arrotondata a due, tre uscite rimanenti uno e pooling due. Prova le modifiche separatamente. Più PE o il salto dei pesi zero richiedono una nuova pianificazione esplicita conservando i valori.

File eseguibile

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));

Progetto finale RTL: convoluzione intera

Gli allegati cnn_window3x3.sv e cnn_conv3x3.sv collegano le finestre a nove prodotti 8×8 con segno, una somma/bias a 24 bit con segno, divisione per due con arrotondamento al pari, saturazione int8, ReLU e un massimo 2×2. La geometria segue il precedente contratto 4×4, a un solo canale, 3×3 senza riempimento. Le scale di ingresso/pesi sono uno, la scala di uscita due e tutti i punti zero sono zero. Il kernel predefinito è [1,0,−1;1,0,−1;1,0,−1], bias otto. I parametri del kernel possono cambiare; ricontrolla il limite della somma a 24 bit.

Confronta separatamente out_acc, out_code e out_activation, con out_row/out_col e trasferimenti di uscita accettati. Done sale dopo l'accettazione del quarto risultato; pooled è allora valido. Sedici ingressi accettati non implicano completamento. Questo RTL usa nove prodotti paralleli e due stadi registrati, anziché la precedente stima a due PE e 24 cicli. Frequenza/potenza richiedono prove di sintesi o del dispositivo.

Esegui il testbench autoverificante condiviso:

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

Gli accumulatori predefiniti sono due, codici/attivazioni uno e uscita del pooling uno. Le verifiche includono limiti con segno, saturazione, equidistanze positive/negative, un kernel denso con segno, pause e reset. Verifica PASS frames=6 windows=24 outputs=24 kernels=2 nel log effettivo. Cambia il primo ingresso a due: la prima somma tre si arrotonda al codice due e pooled diventa due. Poi scegli un altro kernel con segno, calcola a mano la prima finestra, aggiorna la funzione indipendente dei coefficienti e riesegui ogni caso. Non copiare i risultati DUT nel modello di riferimento.