"""Renders the normal-CDF tile from a declared operating point.
Generated for the same reason the exp tile is: its per-segment Taylor
coefficients are transcendental constants at the configured internal
precision, and SystemVerilog cannot compute them at elaboration.
The emitted SystemVerilog mirrors
:meth:`dau_core.fixed_point.NormalCdfDatapath.evaluate` step for step, so
the cocotb bench compares against it with zero tolerance like every other
tile. That model is correctly rounded over EVERY representable input, so for
this tile the algorithm and the mathematics coincide.
Pipelined one multiply per stage, which is what the exp measurements forced:
a DSP48's clock-to-out is ~3.5 ns, so a carry chain sitting between two of
them overruns 8 ns.
"""
# NOTE: no `from __future__ import annotations` — ccflow's Flow.call
# inspects the real annotation objects on __call__
from ccflow import Flow
from .fixed_point import NormalCdfDatapath
from .generators import CoreGenerator, RenderContext, RenderedCore
__all__ = ("NormalCdfTileGenerator",)
def _hex_literal(value: int, width: int) -> str:
return f"{width}'h{value & ((1 << width) - 1):0{(width + 3) // 4}x}"
[docs]
class NormalCdfTileGenerator(CoreGenerator):
"""Renders ``dau_normal_cdf`` for one operating point."""
module: str = "dau_normal_cdf"
datapath: NormalCdfDatapath = NormalCdfDatapath()
@property
def source_name(self) -> str:
return f"{self.module}.sv"
[docs]
@Flow.call
def __call__(self, context: RenderContext) -> RenderedCore: # noqa: ARG002 (ccflow interface)
point = self.datapath
constants = point.constants()
q = point.internal_bits
segments = constants["segments"]
limit = constants["limit"]
wide = q + 3 # coefficients are signed and the constant term reaches 1.0
# one flat coefficient ROM, segment-major: index*(DEGREE+1) + term
flat = [value for segment in constants["coefficients"] for value in segment]
rom = ",\n".join(f" {_hex_literal(value, wide)}" for value in flat)
# Horner, one multiply per stage. Stage 3+2i registers the product and
# 4+2i folds in the next coefficient.
stages = []
for step in range(point.degree):
multiply = 3 + 2 * step
accumulate = multiply + 1
term = point.degree - 1 - step
stages.append(
f" stage_prod[{multiply}] <= PROD_WIDTH'(stage_acc[{multiply - 1}]) * PROD_WIDTH'(stage_res[{multiply - 1}]);"
)
stages.append(
f" stage_acc[{accumulate}] <= WIDE'(WIDE'((stage_prod[{multiply}]) >>> INTERNAL_BITS)"
f" + COEFFICIENT[(stage_seg[{multiply}] * (DEGREE + 1)) + {term}]);"
)
depth = 3 + 2 * point.degree + 1
scale_source = depth - 2
horner_body = "\n".join(stages)
text = f"""`default_nettype none
// GENERATED by dau_core.normal_cdf_generator.NormalCdfTileGenerator -- do
// not edit.
//
// Standard normal CDF over 16-byte quad rows, rendered for ONE operating
// point: internal bits {q}, {segments} segments, degree {point.degree}.
//
// Generated because the per-segment Taylor coefficients are transcendental
// constants at {q} fractional bits, which SystemVerilog cannot compute at
// elaboration. The registry entry names the generator and the operating
// point; editing this file loses the link to the parameters behind it.
//
// Mirrors dau_core.fixed_point.NormalCdfDatapath.evaluate step for step, so
// the bench compares with ZERO TOLERANCE. That model is correctly rounded
// over every one of the 573,441 representable inputs, so here the algorithm
// and the mathematics coincide.
//
// TOTAL: every input has an answer and the tile has NO error path. Beyond
// the domain the result saturates to 0 or 1.0, and that saturation IS the
// correctly rounded answer rather than a clamp. Symmetry is folded, so only
// the positive half is approximated.
module {self.module} #(
parameter int INTERNAL_BITS = {q},
/* verilator lint_off UNUSEDPARAM */
// states the operating point in the source; the segment count it selects
// is already baked into SEGMENTS and the coefficient ROM
parameter int TABLE_BITS = {point.table_bits},
/* verilator lint_on UNUSEDPARAM */
parameter int DEGREE = {point.degree},
parameter int FRACTIONAL_BITS = {point.fractional_bits}
) (
input wire logic clk,
input wire logic rst,
input wire logic [1:0] cfg_dst_field,
input wire logic [1:0] cfg_src_field,
input wire logic input_valid,
output logic input_ready,
input wire logic [63:0] input_data,
input wire logic input_last,
output logic output_valid,
input wire logic output_ready,
output logic [63:0] output_data,
output logic output_last,
output logic status_valid,
input wire logic status_ready,
output logic status_error,
output logic [7:0] status_error_code
);
localparam int WIDE = {wide};
localparam int PROD_WIDTH = 2 * WIDE;
localparam int SEGMENTS = {segments};
localparam int SEG_BITS = {max(point.table_bits, 1)};
localparam int SEGMENT_SHIFT = {constants["segment_shift"]};
localparam logic signed [31:0] LIMIT = 32'sd{limit};
localparam logic [31:0] ONE = 32'd{1 << point.fractional_bits};
localparam logic signed [WIDE-1:0] COEFFICIENT [SEGMENTS * (DEGREE + 1)] = '{{
{rom}
}};
localparam logic [7:0] ERR_STREAM = 8'd2;
localparam int DEPTH = {depth};
typedef enum logic [1:0] {{ IN_W0, IN_W1 }} in_state_t;
typedef enum logic [1:0] {{ OUT_IDLE, OUT_W0, OUT_W1 }} out_state_t;
in_state_t in_state;
out_state_t out_state;
logic batch_open;
logic [1:0] active_dst;
logic [1:0] active_src;
logic status_pending;
logic [7:0] pending_error;
logic [31:0] held_field0;
logic [31:0] held_field1;
logic stage_valid [DEPTH];
logic signed [WIDE-1:0] stage_acc [DEPTH];
logic signed [WIDE-1:0] stage_res [DEPTH];
logic signed [PROD_WIDTH-1:0] stage_prod [DEPTH];
logic [SEG_BITS-1:0] stage_seg [DEPTH];
logic stage_neg [DEPTH];
logic stage_sat [DEPTH];
logic [31:0] stage_field [DEPTH][4];
logic [1:0] stage_dst [DEPTH];
logic stage_last [DEPTH];
logic [31:0] emit_field [4];
logic emit_last;
logic input_fire;
logic output_fire;
logic retire_valid;
logic pipe_advance;
logic pipe_push;
logic [31:0] in_field [4];
logic signed [31:0] src_value;
logic [31:0] magnitude;
logic saturates;
assign in_field[0] = held_field0;
assign in_field[1] = held_field1;
assign in_field[2] = input_data[31:0];
assign in_field[3] = input_data[63:32];
assign src_value = signed'(in_field[active_src]);
assign magnitude = src_value[31] ? (~src_value + 32'sd1) : src_value;
assign saturates = (magnitude >= LIMIT);
assign input_fire = input_valid && input_ready;
assign output_fire = output_valid && output_ready;
assign retire_valid = stage_valid[DEPTH-1];
assign pipe_advance = !retire_valid || (out_state == OUT_IDLE);
assign input_ready = !status_pending && ((in_state == IN_W0) || ((in_state == IN_W1) && pipe_advance));
assign pipe_push = (in_state == IN_W1) && input_fire;
assign output_valid = (out_state == OUT_W0) || (out_state == OUT_W1);
assign output_data = (out_state == OUT_W0) ? {{emit_field[1], emit_field[0]}} : {{emit_field[3], emit_field[2]}};
assign output_last = (out_state == OUT_W1) && emit_last;
assign status_valid = status_pending;
assign status_error = 1'b1;
assign status_error_code = pending_error;
// The segmented range is a POWER OF TWO, so the index is a SHIFT and the
// start a mask. Segmenting the true domain edge instead made both a
// division by a non-power-of-two constant -- free in the Python model,
// a runtime divider in fabric, measured at WNS -9.576 ns.
logic [SEG_BITS-1:0] segment_index;
logic [31:0] segment_start;
assign segment_index = SEG_BITS'(magnitude >> SEGMENT_SHIFT);
assign segment_start = 32'(segment_index) << SEGMENT_SHIFT;
// the rounded, symmetry-folded result
logic signed [WIDE-1:0] rounded;
logic [31:0] clamped;
assign rounded = WIDE'((stage_acc[{scale_source}] + (WIDE'(signed'(1)) <<< (INTERNAL_BITS - FRACTIONAL_BITS - 1))) >>> (INTERNAL_BITS - FRACTIONAL_BITS));
assign clamped = (rounded < 0) ? 32'd0 : ((rounded > WIDE'(signed'({{1'b0, ONE[30:0]}}))) ? ONE : 32'(rounded));
always_ff @(posedge clk) begin
if (rst) begin
in_state <= IN_W0;
out_state <= OUT_IDLE;
batch_open <= 1'b0;
status_pending <= 1'b0;
pending_error <= 8'd0;
emit_last <= 1'b0;
for (int unsigned i = 0; i < DEPTH; i++) stage_valid[i] <= 1'b0;
end else begin
if (pipe_advance) begin
stage_valid[0] <= pipe_push;
if (pipe_push) begin
stage_last[0] <= input_last;
stage_dst[0] <= active_dst;
stage_neg[0] <= src_value[31];
stage_sat[0] <= saturates;
stage_seg[0] <= segment_index;
stage_res[0] <= WIDE'(signed'({{1'b0, magnitude - segment_start}}) <<< (INTERNAL_BITS - FRACTIONAL_BITS));
stage_acc[0] <= COEFFICIENT[(segment_index * (DEGREE + 1)) + DEGREE];
for (int unsigned f = 0; f < 4; f++) stage_field[0][f] <= in_field[f];
end
for (int unsigned i = 1; i < DEPTH; i++) begin
stage_valid[i] <= stage_valid[i-1];
stage_last[i] <= stage_last[i-1];
stage_dst[i] <= stage_dst[i-1];
stage_neg[i] <= stage_neg[i-1];
stage_sat[i] <= stage_sat[i-1];
stage_seg[i] <= stage_seg[i-1];
stage_acc[i] <= stage_acc[i-1];
stage_res[i] <= stage_res[i-1];
stage_prod[i] <= stage_prod[i-1];
for (int unsigned f = 0; f < 4; f++) stage_field[i][f] <= stage_field[i-1][f];
end
{horner_body}
end
case (in_state)
IN_W0:
if (input_fire) begin
if (input_last) begin
pending_error <= ERR_STREAM;
status_pending <= 1'b1;
end else begin
held_field0 <= input_data[31:0];
held_field1 <= input_data[63:32];
if (!batch_open) begin
batch_open <= 1'b1;
active_dst <= cfg_dst_field;
active_src <= cfg_src_field;
end
in_state <= IN_W1;
end
end
IN_W1:
if (input_fire) begin
in_state <= IN_W0;
if (input_last) batch_open <= 1'b0;
end
default: in_state <= IN_W0;
endcase
case (out_state)
OUT_IDLE:
if (retire_valid) begin
automatic logic [31:0] value;
value = stage_sat[DEPTH-1] ? (stage_neg[DEPTH-1] ? 32'd0 : ONE)
: (stage_neg[DEPTH-1] ? (ONE - clamped) : clamped);
for (int unsigned f = 0; f < 4; f++) emit_field[f] <= stage_field[DEPTH-1][f];
emit_field[stage_dst[DEPTH-1]] <= value;
emit_last <= stage_last[DEPTH-1];
out_state <= OUT_W0;
end
OUT_W0: if (output_fire) out_state <= OUT_W1;
OUT_W1: if (output_fire) out_state <= OUT_IDLE;
default: out_state <= OUT_IDLE;
endcase
if (status_pending && status_ready) begin
status_pending <= 1'b0;
in_state <= IN_W0;
batch_open <= 1'b0;
for (int unsigned i = 0; i < DEPTH; i++) stage_valid[i] <= 1'b0;
end
end
end
endmodule
`default_nettype wire
"""
return RenderedCore(
module=self.module,
source_name=self.source_name,
text=text,
parameters={
"INTERNAL_BITS": q,
"TABLE_BITS": point.table_bits,
"DEGREE": point.degree,
"FRACTIONAL_BITS": point.fractional_bits,
},
)