캡스톤: 작은 추론 경로와 보고서

목표와 제출 기준

앞 실습의 model.mjs와 같은 폴더에 capstone-lab.mjs를 만들고 아래 코드를 저장한 뒤 node capstone-lab.mjs를 실행합니다. 목표는 한 층의 합성곱,바이어스,재양자화,ReLU와 최대 풀링을 연결하고 값과 비용을 분리해 보고하는 것입니다. 전체 분류 모델이나 학습기를 만들었다는 목표가 아닙니다. 설치된 Digital Design Studio가 지원하는 파일 생성·열기 흐름에서는 같은 파일을 수동으로 검토할 수 있으며 이 자료는 DDS 안에서의 실행 완료를 의미하지 않습니다.

완료 기준은 직접 모델과 펼친 패치 내적의 누적 일치,중간 출력·활성화·풀링의 기대값 일치,모양과 MAC 수,명시적 일정의 사이클 수와 슬롯 사용률,유일한 텐서 바이트를 기록하는 것입니다. 추가로 입력 한 값과 커널 한 값을 바꾼 새 손계산 fixture를 제출하세요. 실행 로그에는 명령,파일 버전과 가정을 함께 남깁니다. 비교 실패를 기대값만 고쳐 숨기지 말고 첫 오차 항을 추적해야 합니다.

기준 결과 계산

입력은 4×4의 1부터 16이고 커널은 [1,0,−1;1,0,−1;1,0,−1]입니다. 각 3×3 창은 세 행마다 왼쪽에서 오른쪽을 빼므로 −2를 세 번 더한 −6입니다. 바이어스 8을 더하면 네 위치 모두 누적 2입니다. 출력 multiplier 0.5,z=0은 코드 1을 만들고 ReLU도 1을 유지합니다. 2×2 최대 풀링의 최종 결과는 [1]입니다. 원소 값이 같아도 창 위치와 합계 항을 확인해야 하므로 이전의 비대칭 경계 fixture도 같이 유지합니다.

출력 모양은 [M,Ho,Wo]=[1,2,2]입니다. 총 MAC은 4출력×9항=36이며 0가중치 항도 일정에 포함합니다. 두 단일-MAC PE가 출력 하나씩 감소 항을 나눠 처리한다고 가정하면 ceil(9/2)=5MAC사이클과 출력당 추가 1사이클입니다. 총 모델 사이클은 4×6=24이고 MAC 구간 슬롯 사용률은 36/(4×5×2)=90%입니다. 이 수는 실제 장치에서 측정한 지연이 아닙니다.

저장과 전송 보고

유일한 int8 입력은 16바이트,가중치 9바이트,합성곱 출력은 4바이트이고 32비트 바이어스는 4바이트입니다. 이는 원본 텐서 크기이며 실제 전송량에는 중첩 창의 재읽기,부분합과 제어 비용이 추가될 수 있습니다. 풀링 후 출력은 1바이트입니다. 입력의 정수 scale은 1이고 가중치 scale은 1,출력 scale은 2로 두어 multiplier 0.5를 해석할 수 있습니다. 코드와 배율을 함께 기록하세요.

확장 과제

바이어스를 0으로 바꾸면 누적 −6,재양자화 −3,ReLU 0과 풀링 0을 예상합니다. 입력 첫 값 1을 2로 바꾸면 첫 누적만 3이 되어 1.5의 짝수 반올림이 2를 만듭니다. 다른 세 출력은 1이고 최종 풀링은 2입니다. 두 변경을 한꺼번에 하기 전에 각각 실행하고 기대값을 작성합니다. 추가 PE나 0가중치 생략 최적화는 결과를 유지하면서 가정과 일정 수를 새로 검증해야 합니다.

실행 코드

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는 창 생성기 뒤에 9개의 signed 8×8 곱,24비트 signed 누적,바이어스,짝수 반올림의 2 나누기,int8 포화,ReLU와 2×2 최대값을 연결합니다. 형상은 앞 RTL 실습과 같은 4×4·한 채널·3×3 valid입니다. scale은 입력·가중치 1,출력 2이고 모든 zero point는 0입니다. 가중치는 [1,0,−1;1,0,−1;1,0,−1],바이어스는 8입니다. 가중치 파라미터를 바꿀 수 있지만 누적과 바이어스가 24비트 범위에 들어가는지 다시 계산하세요.

out_acc,out_code,out_activation은 각각 중간 누적,포화된 출력 코드와 ReLU 결과입니다. out_row/out_col로 네 출력의 순서를 비교하고 out_valid·out_ready의 수락만 셉니다. 마지막 출력이 수락된 뒤 done이 올라가며 pooled가 유효합니다. 입력 16개를 받았다는 사실만으로 done을 선언하지 않습니다. 이 RTL은 9개 곱을 병렬로 계산하는 두 registered stage이며 앞의 2PE·24사이클 소프트웨어 일정과 다른 구현입니다. 합성 주파수나 전력은 이 자료로 확인할 수 없습니다.

같은 폴더의 tb_cnn_stream.sv로 실행하세요.

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

기본 입력의 네 누적은 2,코드·ReLU는 1,pooled는 1입니다. 검사에는 음수·양수 절반,부호 경계값,포화,밀집 signed 커널,멈춤과 reset이 포함됩니다. PASS frames=6 windows=24 outputs=24 kernels=2를 실제 로그에서 확인하세요. 변경 과제: 첫 입력을 2로 바꾸면 첫 누적 3이 코드 2로 반올림되고 pooled=2입니다. 다른 signed 커널을 선택해 첫 창을 손계산하고 testbench의 독립 coefficient 함수를 바꾼 뒤 모든 fixture를 다시 실행하세요. DUT 결과를 복사해 기대값을 만들지 마세요.