def leaf_cost

stelliferous@0.0.2.0 · kernel_planning.bend:135

import Base
import stelliferous@0.0.2.0/kernel_cost.bend as Cost
import stelliferous@0.0.2.0/kernel_pipeline.bend as Pipeline

135def leaf_cost(span: U32,+count: U32,work: Work,+config: Cost.Configuration) -> Movement:
136  Work{positions,+inner,+rows,owners,geometry,activation} = work
137  Cost.Configuration{mac_ns,+byte_ns,operation_ns,grain,wake} = config
138  +span = span
139  +packed = Nat.mul(32n,capacity(U32.max(Pipeline.block_tiles(),inner)))
140  +output_capacity = capacity((rows * count : U32))
141  +output = Nat.mul(4n,Nat.max(2n,output_capacity))
142  # Padded lanes still execute every MAC, including absent rows of the last unit.
143  +core = (U32.to_f32(((rows + 3) / 4 * 4 : U32)) * U32.to_f32(span) * U32.to_f32(inner) * mac_ns : F32)
144  # Initializing each allocation and transferring logical data are separate work.
145  +packing = ((U32.to_f32(span) * U32.to_f32(inner) * 4.0 + payload_float(packed)) * byte_ns : F32)
146  +stores = ((U32.to_f32(rows) * U32.to_f32(count) * 4.0 + payload_float(Nat.mul(4n,output_capacity))) * byte_ns : F32)
147  +finish = (U32.to_f32(rows) * U32.to_f32(count) * activation_ns(activation,config) : F32)
148  +time = (core + packing + stores + finish : F32)
149  Movement{packed,output,0n,time,time,time}