KernelPlanning.leaf_cost: def of stelliferous@0.0.2.0
Import: import stelliferous@0.0.2.0/kernel_planning.bend as KernelPlanning
Lines 135-149 of stelliferous@0.0.2.0/kernel_planning.bend: https://jend.leeeet.dev/read.txt?path=stelliferous@0.0.2.0/kernel_planning.bend&start_line=135&end_line=149

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

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

Uses (18):
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=F32
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=F32.add
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=F32.mul
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=Nat.max
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=Nat.mul
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32.add
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32.div
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32.max
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32.mul
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32.to_f32
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_cost.bend&name=Configuration
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_pipeline.bend&name=block_tiles
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_planning.bend&name=Movement
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_planning.bend&name=Work
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_planning.bend&name=activation_ns
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_planning.bend&name=capacity
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_planning.bend&name=payload_float