KernelPipeline.blocked: def of stelliferous@0.0.2.0
Import: import stelliferous@0.0.2.0/kernel_pipeline.bend as KernelPipeline
Lines 151-159 of stelliferous@0.0.2.0/kernel_pipeline.bend: https://jend.leeeet.dev/read.txt?path=stelliferous@0.0.2.0/kernel_pipeline.bend&start_line=151&end_line=159

import Base
import stelliferous@0.0.2.0/kernel_tile.bend as Tile
import stelliferous@0.0.2.0/storage_buffer.bend as Storage
import stelliferous@0.0.2.0/tensor.bend as Tensors
import stelliferous@0.0.2.0/kernel_window.bend as KernelWindow
import stelliferous@0.0.2.0/math_activation.bend as Activations

# Exactly one reusable tile buffer and one output buffer per leaf.
# An inner larger than the target still needs one complete eight-position group.
def blocked(~custom: F32 -> F32,+activation: Activations.Activation,+work: Work,pack_start: U32,input: Storage.Buffer<F32>,weights: Storage.Buffer<Tile.Tile8>,bias: Storage.Buffer<F32>) -> Storage.Buffer<F32>:
  Work{+rows,h,w,k,s,p,ow,+positions,+inner} = work
  +padded = KernelWindow.padded(positions)
  +span = block_positions(inner)
  blocks_finished(~custom,activation,(rows * positions : U32),
    blocks(U32.to_nat(((padded + span - 1) / span : U32)),work,pack_start,0,padded,span,
      BlockState{input,Tensors.buffer(~Tile.Tile8,Tile.filled(0.0),U32.max(block_tiles(),inner)),weights,bias,Tensors.buffer(~F32,0.0,(rows * positions : U32))}))

Uses (20):
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=F32
  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.sub
  https://jend.leeeet.dev/definition.txt?path=Base/base.bend&name=U32.to_nat
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_pipeline.bend&name=BlockState
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_pipeline.bend&name=Work
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_pipeline.bend&name=block_positions
  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_pipeline.bend&name=blocks
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_pipeline.bend&name=blocks_finished
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_tile.bend&name=Tile8
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_tile.bend&name=filled
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/kernel_window.bend&name=padded
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/math_activation.bend&name=Activation
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/storage_buffer.bend&name=Buffer
  https://jend.leeeet.dev/definition.txt?path=stelliferous@0.0.2.0/tensor.bend&name=buffer