Write compute programs
in the language you think in.

incipit {
    fixum lista<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    fixum lista<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    fixum tf32[] seed ← vacua
    fixum tf32[2, 3] a ← seed.strue(flat_a, [2, 3])
    fixum tf32[3, 4] b ← seed.strue(flat_b, [3, 4])
    fixum tf32[2, 4] product ← a.matmul(b)
    fixum f32 mean ← product.media()
    nota mean
}

เริ่ม {
    คงที่ รายการ<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    คงที่ รายการ<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    คงที่ เทนเซอร์<f32, []> seed ← เซตว่าง
    คงที่ เทนเซอร์<f32, [2, 3]> a ← seed.strue(flat_a, [2, 3])
    คงที่ เทนเซอร์<f32, [3, 4]> b ← seed.strue(flat_b, [3, 4])
    คงที่ เทนเซอร์<f32, [2, 4]> product ← a.matmul(b)
    คงที่ f32 mean ← product.media()
    บันทึก mean
}

入口 {
    常量 列表<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    常量 列表<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    常量 张量<f32, []> seed ← 空集
    常量 张量<f32, [2, 3]> a ← seed.strue(flat_a, [2, 3])
    常量 张量<f32, [3, 4]> b ← seed.strue(flat_b, [3, 4])
    常量 张量<f32, [2, 4]> product ← a.matmul(b)
    常量 f32 mean ← product.media()
    显示 mean
}

入口 {
    定值 列表<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    定值 列表<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    定值 張量<f32, []> seed ← 空集
    定值 張量<f32, [2, 3]> a ← seed.strue(flat_a, [2, 3])
    定值 張量<f32, [3, 4]> b ← seed.strue(flat_b, [3, 4])
    定值 張量<f32, [2, 4]> product ← a.matmul(b)
    定值 f32 mean ← product.media()
    註記 mean
}

bắt_đầu {
    hằng danh_sách<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    hằng danh_sách<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    hằng ten_xo<f32, []> seed ← tập_rỗng
    hằng ten_xo<f32, [2, 3]> a ← seed.strue(flat_a, [2, 3])
    hằng ten_xo<f32, [3, 4]> b ← seed.strue(flat_b, [3, 4])
    hằng ten_xo<f32, [2, 4]> product ← a.matmul(b)
    hằng f32 mean ← product.media()
    ghi_chú mean
}

بداية {
    ثابت قائمة<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    ثابت قائمة<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    ثابت موتر<f32, []> seed ← فارغ
    ثابت موتر<f32, [2, 3]> a ← seed.strue(flat_a, [2, 3])
    ثابت موتر<f32, [3, 4]> b ← seed.strue(flat_b, [3, 4])
    ثابت موتر<f32, [2, 4]> product ← a.matmul(b)
    ثابت f32 mean ← product.media()
    اعرض mean
}

आरंभ {
    स्थिर सूची<f32> flat_a ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]
    स्थिर सूची<f32> flat_b ← [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0]
    स्थिर टेंसर<f32, []> seed ← खाली
    स्थिर टेंसर<f32, [2, 3]> a ← seed.strue(flat_a, [2, 3])
    स्थिर टेंसर<f32, [3, 4]> b ← seed.strue(flat_b, [3, 4])
    स्थिर टेंसर<f32, [2, 4]> product ← a.matmul(b)
    स्थिर f32 mean ← product.media()
    दिखाओ mean
}

// Generated by radix - do not edit

package main

import (
    "fmt"
    "strconv"
    "strings"
)

type faberTensor[T any] struct {
    data []T
    shape []int
}

func faberTensorElementCount(shape []int) int {
    const maxInt = int(^uint(0) >> 1)
    total := 1
    for _, dim := range shape {
        if dim < 0 { panic("tensor shape dimension must be non-negative") }
        if dim > 0 && total > maxInt/dim { panic("tensor shape element count overflow") }
        total *= dim
    }
    return total
}

func faberIndexSlice(indices any) []int {
    switch values := indices.(type) {
        case []int:
            return append([]int{}, values...)
        case []uint32:
            out := make([]int, len(values)); for i, value := range values { out[i] = int(value) }; return out
        case []uint64:
            out := make([]int, len(values)); for i, value := range values { out[i] = int(value) }; return out
        case []int32:
            out := make([]int, len(values)); for i, value := range values { out[i] = int(value) }; return out
        case []int64:
            out := make([]int, len(values)); for i, value := range values { out[i] = int(value) }; return out
        default:
            panic("tensor index must be a numeric list")
    }
}

func faberTensorOffset(shape []int, rawIndices any) *int {
    const maxInt = int(^uint(0) >> 1)
    indices := faberIndexSlice(rawIndices)
    if len(indices) != len(shape) { return nil }
    offset := 0
    stride := 1
    for axis := len(shape) - 1; axis >= 0; axis-- {
        idx := indices[axis]
        dim := shape[axis]
        if dim < 0 || idx < 0 || idx >= dim { return nil }
        if idx > 0 && stride > (maxInt-offset)/idx { return nil }
        offset += idx * stride
        if dim > 0 && stride > maxInt/dim { return nil }
        stride *= dim
    }
    return &offset
}

func (t faberTensor[T]) Crea(fill T, shape []int) faberTensor[T] {
    data := make([]T, faberTensorElementCount(shape))
    for i := range data { data[i] = fill }
    return faberTensor[T]{data: data, shape: append([]int{}, shape...)}
}

func (t faberTensor[T]) Strue(data []T, shape []int) faberTensor[T] {
    if faberTensorElementCount(shape) != len(data) { panic("tensor structa element count does not match shape") }
    return faberTensor[T]{data: append([]T{}, data...), shape: append([]int{}, shape...)}
}

func (t faberTensor[T]) Longitudo() int { return len(t.shape) }
func (t faberTensor[T]) Magnitudines() []int { return append([]int{}, t.shape...) }
func (t faberTensor[T]) Planata() []T { return append([]T{}, t.data...) }
func (t faberTensor[T]) Materialize() faberTensor[T] { return faberTensor[T]{data: append([]T{}, t.data...), shape: append([]int{}, t.shape...)} }

func faberTensorAdd[T any](left T, right T) T {
    switch value := any(left).(type) {
        case int: return any(value + any(right).(int)).(T)
        case int32: return any(value + any(right).(int32)).(T)
        case int64: return any(value + any(right).(int64)).(T)
        case uint: return any(value + any(right).(uint)).(T)
        case uint32: return any(value + any(right).(uint32)).(T)
        case uint64: return any(value + any(right).(uint64)).(T)
        case float32: return any(value + any(right).(float32)).(T)
        case float64: return any(value + any(right).(float64)).(T)
        default: panic("tensor arithmetic requires numeric elements")
    }
}

func faberTensorMul[T any](left T, right T) T {
    switch value := any(left).(type) {
        case int: return any(value * any(right).(int)).(T)
        case int32: return any(value * any(right).(int32)).(T)
        case int64: return any(value * any(right).(int64)).(T)
        case uint: return any(value * any(right).(uint)).(T)
        case uint32: return any(value * any(right).(uint32)).(T)
        case uint64: return any(value * any(right).(uint64)).(T)
        case float32: return any(value * any(right).(float32)).(T)
        case float64: return any(value * any(right).(float64)).(T)
        default: panic("tensor arithmetic requires numeric elements")
    }
}

func faberTensorSub[T any](left T, right T) T {
    switch value := any(left).(type) {
        case int: return any(value - any(right).(int)).(T)
        case int32: return any(value - any(right).(int32)).(T)
        case int64: return any(value - any(right).(int64)).(T)
        case uint: return any(value - any(right).(uint)).(T)
        case uint32: return any(value - any(right).(uint32)).(T)
        case uint64: return any(value - any(right).(uint64)).(T)
        case float32: return any(value - any(right).(float32)).(T)
        case float64: return any(value - any(right).(float64)).(T)
        default: panic("tensor arithmetic requires numeric elements")
    }
}

func faberTensorShapeEqual(left []int, right []int) bool {
    if len(left) != len(right) { return false }
    for i, dim := range left { if dim != right[i] { return false } }
    return true
}

func faberTensorMean[T any](data []T) T {
    if len(data) == 0 { panic("tensor media requires non-empty data") }
    switch any(data[0]).(type) {
        case float32:
            var total float32
            for _, value := range data { total += any(value).(float32) }
            return any(total / float32(len(data))).(T)
        case float64:
            var total float64
            for _, value := range data { total += any(value).(float64) }
            return any(total / float64(len(data))).(T)
        default: panic("tensor media requires floating-point elements")
    }
}

func (t faberTensor[T]) Summa() T {
    var total T
    for _, value := range t.data { total = faberTensorAdd(total, value) }
    return total
}

func (t faberTensor[T]) Media() T { return faberTensorMean(t.data) }

func (a faberTensor[T]) Addita(b faberTensor[T]) faberTensor[T] {
    if !faberTensorShapeEqual(a.shape, b.shape) { panic("tensor elementwise arithmetic requires equal shapes") }
    data := make([]T, len(a.data))
    for i := range data { data[i] = faberTensorAdd(a.data[i], b.data[i]) }
    return faberTensor[T]{data: data, shape: append([]int{}, a.shape...)}
}

func (a faberTensor[T]) Subtrahe(b faberTensor[T]) faberTensor[T] {
    if !faberTensorShapeEqual(a.shape, b.shape) { panic("tensor elementwise arithmetic requires equal shapes") }
    data := make([]T, len(a.data))
    for i := range data { data[i] = faberTensorSub(a.data[i], b.data[i]) }
    return faberTensor[T]{data: data, shape: append([]int{}, a.shape...)}
}

func (a faberTensor[T]) Multiplica(b faberTensor[T]) faberTensor[T] {
    if !faberTensorShapeEqual(a.shape, b.shape) { panic("tensor elementwise arithmetic requires equal shapes") }
    data := make([]T, len(a.data))
    for i := range data { data[i] = faberTensorMul(a.data[i], b.data[i]) }
    return faberTensor[T]{data: data, shape: append([]int{}, a.shape...)}
}

func (a faberTensor[T]) Matmul(b faberTensor[T]) faberTensor[T] {
    if len(a.shape) != 2 || len(b.shape) != 2 || a.shape[1] != b.shape[0] { panic("tensor matmul requires compatible rank-2 shapes") }
    rows, inner, cols := a.shape[0], a.shape[1], b.shape[1]
    data := make([]T, rows*cols)
    for row := 0; row < rows; row++ {
        for col := 0; col < cols; col++ {
            var sum T
            for k := 0; k < inner; k++ { sum = faberTensorAdd(sum, faberTensorMul(a.data[row*inner+k], b.data[k*cols+col])) }
            data[row*cols+col] = sum
        }
    }
    return faberTensor[T]{data: data, shape: []int{rows, cols}}
}

func (t faberTensor[T]) Forma(shape []int) faberTensor[T] {
    if faberTensorElementCount(shape) != len(t.data) { panic("tensor forma (reshape) element count mismatch") }
    return faberTensor[T]{data: append([]T{}, t.data...), shape: append([]int{}, shape...)}
}

func (t faberTensor[T]) Accipe(indices any) *T {
    offset := faberTensorOffset(t.shape, indices)
    if offset == nil || *offset < 0 || *offset >= len(t.data) { return nil }
    return &t.data[*offset]
}

func (t *faberTensor[T]) Ponde(indices any, value T) {
    offset := faberTensorOffset(t.shape, indices)
    if offset == nil || *offset < 0 || *offset >= len(t.data) { panic("tensor ponde invalid index") }
    t.data[*offset] = value
}

func (t *faberTensor[T]) Reple(value T) {
    for i := range t.data { t.data[i] = value }
}

func (t faberTensor[T]) Sectio(start int, end int) faberTensor[T] {
    if len(t.shape) == 0 || start < 0 || end < start || end > t.shape[0] { panic("tensor sectio invalid slice bounds") }
    inner := faberTensorElementCount(t.shape[1:])
    shape := append([]int{end - start}, t.shape[1:]...)
    return faberTensor[T]{data: append([]T{}, t.data[start*inner:end*inner]...), shape: shape}
}

func main() {
    flat_a := []float32{float32(1), float32(2), float32(3), float32(4), float32(5), float32(6)}
    flat_b := []float32{float32(1), float32(2), float32(3), float32(4), float32(5), float32(6), float32(7), float32(8), float32(9), float32(10), float32(11), float32(12)}
    seed := faberTensor[float32]{shape: []int{}}
    a := seed.Strue(flat_a, []int{2, 3})
    b := seed.Strue(flat_b, []int{3, 4})
    product := a.Matmul(b)
    mean := float32(product.Media())
    fmt.Println(func(v float64) string { s := strconv.FormatFloat(v, 'f', -1, 64); if !strings.ContainsAny(s, ".eE") { return s + ".0" }; return s }(float64(mean)))
}

// … 285 lines of generated display/runtime shim elided …

        const flat_a: Array<number> = [1, 2, 3, 4, 5, 6];
        const flat_b: Array<number> = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12];
        const seed: FaberTensor<number> = FaberTensor.empty<number>([]);
        const a: FaberTensor<number> = seed.strue(flat_a, [2, 3]);
        const b: FaberTensor<number> = seed.strue(flat_b, [3, 4]);
        const product: FaberTensor<number> = a.matmul(b);
        const mean: number = product.media();
        console.log(__faberDisplay(mean, "fractus"));
    }})();

; Generated by radix MIR LLVM IR probe - experimental artifact.
%FaberRtSliceV1 = type { ptr, i64 }
%FaberRtExitV1 = type i64
%FaberRtPtrResultV1 = type { i32, ptr }
%FaberRtStatusV1 = type { i32 }
@__faber_rt_v1_context = linkonce_odr global ptr null
@__faber_rt_v1_status = linkonce_odr global i32 0
declare void @__faber_rt_v1_fatal(ptr, %FaberRtSliceV1) noreturn
declare void @__faber_rt_v1_numerus_overflow(ptr) noreturn


; @runtime __faber_rt_v1_array_new category=core-semantics
declare %FaberRtPtrResultV1 @__faber_rt_v1_array_new(ptr, i32)
; @runtime __faber_rt_v1_array_push category=core-semantics
declare i32 @__faber_rt_v1_array_push(ptr, ptr, i32, ptr)
; @runtime __faber_rt_v1_diagnostic_nota_f32 category=host-integration
declare i32 @__faber_rt_v1_diagnostic_nota_f32(ptr, float)
; @runtime __faber_rt_v1_tensor_from_flat category=core-semantics
declare %FaberRtPtrResultV1 @__faber_rt_v1_tensor_from_flat(ptr, i32, ptr, ptr)
; @runtime __faber_rt_v1_tensor_matmul category=core-semantics
declare %FaberRtPtrResultV1 @__faber_rt_v1_tensor_matmul(ptr, ptr, ptr)
; @runtime __faber_rt_v1_tensor_mean category=core-semantics
declare i32 @__faber_rt_v1_tensor_mean(ptr, ptr, i32, ptr)
; @runtime __faber_rt_v1_tensor_new category=core-semantics
declare %FaberRtPtrResultV1 @__faber_rt_v1_tensor_new(ptr, i32)

define void @incipit() {
    entry:
      %l0.addr = alloca ptr
      %l1.addr = alloca ptr
      %l2.addr = alloca ptr
      %l3.addr = alloca ptr
      %l4.addr = alloca ptr
      %l5.addr = alloca ptr
      %l6.addr = alloca float
      %t0.addr = alloca float
      %t1.addr = alloca float
      %t2.addr = alloca float
      %t3.addr = alloca float
      %t4.addr = alloca float
      %t5.addr = alloca float
      %t6.addr = alloca ptr
      %t7.addr = alloca float
      %t8.addr = alloca float
      %t9.addr = alloca float
      %t10.addr = alloca float
      %t11.addr = alloca float
      %t12.addr = alloca float
      %t13.addr = alloca float
      %t14.addr = alloca float
      %t15.addr = alloca float
      %t16.addr = alloca float
      %t17.addr = alloca float
      %t18.addr = alloca float
      %t19.addr = alloca ptr
      %t20.addr = alloca ptr
      %t21.addr = alloca ptr
      %t22.addr = alloca ptr
      %t23.addr = alloca ptr
      %t24.addr = alloca ptr
      %t25.addr = alloca ptr
      %t26.addr = alloca float
      br label %b0
    b0:
      %cast0 = fptrunc double 1.0 to float
      store float %cast0, ptr %t0.addr
      %cast1 = fptrunc double 2.0 to float
      store float %cast1, ptr %t1.addr
      %cast2 = fptrunc double 3.0 to float
      store float %cast2, ptr %t2.addr
      %cast3 = fptrunc double 4.0 to float
      store float %cast3, ptr %t3.addr
      %cast4 = fptrunc double 5.0 to float
      store float %cast4, ptr %t4.addr
      %cast5 = fptrunc double 6.0 to float
      store float %cast5, ptr %t5.addr
      %faber.context6 = load ptr, ptr @__faber_rt_v1_context
      %faber.array.result6 = call %FaberRtPtrResultV1 @__faber_rt_v1_array_new(ptr %faber.context6, i32 5)
      %faber.array.status6 = extractvalue %FaberRtPtrResultV1 %faber.array.result6, 0
      %faber.array.handle6 = extractvalue %FaberRtPtrResultV1 %faber.array.result6, 1
      %faber.old.status6 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error6 = icmp ne i32 %faber.old.status6, 0
      %faber.latched.status6 = select i1 %faber.has.error6, i32 %faber.old.status6, i32 %faber.array.status6
      store i32 %faber.latched.status6, ptr @__faber_rt_v1_status
      %load8 = load float, ptr %t0.addr
      %faber.array.value7 = alloca float
      store float %load8, ptr %faber.array.value7
      %faber.array.status7 = call i32 @__faber_rt_v1_array_push(ptr %faber.context6, ptr %faber.array.handle6, i32 5, ptr %faber.array.value7)
      %faber.old.status7 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error7 = icmp ne i32 %faber.old.status7, 0
      %faber.latched.status7 = select i1 %faber.has.error7, i32 %faber.old.status7, i32 %faber.array.status7
      store i32 %faber.latched.status7, ptr @__faber_rt_v1_status
      %load10 = load float, ptr %t1.addr
      %faber.array.value9 = alloca float
      store float %load10, ptr %faber.array.value9
      %faber.array.status9 = call i32 @__faber_rt_v1_array_push(ptr %faber.context6, ptr %faber.array.handle6, i32 5, ptr %faber.array.value9)
      %faber.old.status9 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error9 = icmp ne i32 %faber.old.status9, 0
      %faber.latched.status9 = select i1 %faber.has.error9, i32 %faber.old.status9, i32 %faber.array.status9
      store i32 %faber.latched.status9, ptr @__faber_rt_v1_status
      %load12 = load float, ptr %t2.addr
      %faber.array.value11 = alloca float
      store float %load12, ptr %faber.array.value11
      %faber.array.status11 = call i32 @__faber_rt_v1_array_push(ptr %faber.context6, ptr %faber.array.handle6, i32 5, ptr %faber.array.value11)
      %faber.old.status11 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error11 = icmp ne i32 %faber.old.status11, 0
      %faber.latched.status11 = select i1 %faber.has.error11, i32 %faber.old.status11, i32 %faber.array.status11
      store i32 %faber.latched.status11, ptr @__faber_rt_v1_status
      %load14 = load float, ptr %t3.addr
      %faber.array.value13 = alloca float
      store float %load14, ptr %faber.array.value13
      %faber.array.status13 = call i32 @__faber_rt_v1_array_push(ptr %faber.context6, ptr %faber.array.handle6, i32 5, ptr %faber.array.value13)
      %faber.old.status13 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error13 = icmp ne i32 %faber.old.status13, 0
      %faber.latched.status13 = select i1 %faber.has.error13, i32 %faber.old.status13, i32 %faber.array.status13
      store i32 %faber.latched.status13, ptr @__faber_rt_v1_status
      %load16 = load float, ptr %t4.addr
      %faber.array.value15 = alloca float
      store float %load16, ptr %faber.array.value15
      %faber.array.status15 = call i32 @__faber_rt_v1_array_push(ptr %faber.context6, ptr %faber.array.handle6, i32 5, ptr %faber.array.value15)
      %faber.old.status15 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error15 = icmp ne i32 %faber.old.status15, 0
      %faber.latched.status15 = select i1 %faber.has.error15, i32 %faber.old.status15, i32 %faber.array.status15
      store i32 %faber.latched.status15, ptr @__faber_rt_v1_status
      %load18 = load float, ptr %t5.addr
      %faber.array.value17 = alloca float
      store float %load18, ptr %faber.array.value17
      %faber.array.status17 = call i32 @__faber_rt_v1_array_push(ptr %faber.context6, ptr %faber.array.handle6, i32 5, ptr %faber.array.value17)
      %faber.old.status17 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error17 = icmp ne i32 %faber.old.status17, 0
      %faber.latched.status17 = select i1 %faber.has.error17, i32 %faber.old.status17, i32 %faber.array.status17
      store i32 %faber.latched.status17, ptr @__faber_rt_v1_status
      store ptr %faber.array.handle6, ptr %t6.addr
      %load19 = load ptr, ptr %t6.addr
      store ptr %load19, ptr %l0.addr
      %cast20 = fptrunc double 1.0 to float
      store float %cast20, ptr %t7.addr
      %cast21 = fptrunc double 2.0 to float
      store float %cast21, ptr %t8.addr
      %cast22 = fptrunc double 3.0 to float
      store float %cast22, ptr %t9.addr
      %cast23 = fptrunc double 4.0 to float
      store float %cast23, ptr %t10.addr
      %cast24 = fptrunc double 5.0 to float
      store float %cast24, ptr %t11.addr
      %cast25 = fptrunc double 6.0 to float
      store float %cast25, ptr %t12.addr
      %cast26 = fptrunc double 7.0 to float
      store float %cast26, ptr %t13.addr
      %cast27 = fptrunc double 8.0 to float
      store float %cast27, ptr %t14.addr
      %cast28 = fptrunc double 9.0 to float
      store float %cast28, ptr %t15.addr
      %cast29 = fptrunc double 10.0 to float
      store float %cast29, ptr %t16.addr
      %cast30 = fptrunc double 11.0 to float
      store float %cast30, ptr %t17.addr
      %cast31 = fptrunc double 12.0 to float
      store float %cast31, ptr %t18.addr
      %faber.context32 = load ptr, ptr @__faber_rt_v1_context
      %faber.array.result32 = call %FaberRtPtrResultV1 @__faber_rt_v1_array_new(ptr %faber.context32, i32 5)
      %faber.array.status32 = extractvalue %FaberRtPtrResultV1 %faber.array.result32, 0
      %faber.array.handle32 = extractvalue %FaberRtPtrResultV1 %faber.array.result32, 1
      %faber.old.status32 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error32 = icmp ne i32 %faber.old.status32, 0
      %faber.latched.status32 = select i1 %faber.has.error32, i32 %faber.old.status32, i32 %faber.array.status32
      store i32 %faber.latched.status32, ptr @__faber_rt_v1_status
      %load34 = load float, ptr %t7.addr
      %faber.array.value33 = alloca float
      store float %load34, ptr %faber.array.value33
      %faber.array.status33 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value33)
      %faber.old.status33 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error33 = icmp ne i32 %faber.old.status33, 0
      %faber.latched.status33 = select i1 %faber.has.error33, i32 %faber.old.status33, i32 %faber.array.status33
      store i32 %faber.latched.status33, ptr @__faber_rt_v1_status
      %load36 = load float, ptr %t8.addr
      %faber.array.value35 = alloca float
      store float %load36, ptr %faber.array.value35
      %faber.array.status35 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value35)
      %faber.old.status35 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error35 = icmp ne i32 %faber.old.status35, 0
      %faber.latched.status35 = select i1 %faber.has.error35, i32 %faber.old.status35, i32 %faber.array.status35
      store i32 %faber.latched.status35, ptr @__faber_rt_v1_status
      %load38 = load float, ptr %t9.addr
      %faber.array.value37 = alloca float
      store float %load38, ptr %faber.array.value37
      %faber.array.status37 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value37)
      %faber.old.status37 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error37 = icmp ne i32 %faber.old.status37, 0
      %faber.latched.status37 = select i1 %faber.has.error37, i32 %faber.old.status37, i32 %faber.array.status37
      store i32 %faber.latched.status37, ptr @__faber_rt_v1_status
      %load40 = load float, ptr %t10.addr
      %faber.array.value39 = alloca float
      store float %load40, ptr %faber.array.value39
      %faber.array.status39 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value39)
      %faber.old.status39 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error39 = icmp ne i32 %faber.old.status39, 0
      %faber.latched.status39 = select i1 %faber.has.error39, i32 %faber.old.status39, i32 %faber.array.status39
      store i32 %faber.latched.status39, ptr @__faber_rt_v1_status
      %load42 = load float, ptr %t11.addr
      %faber.array.value41 = alloca float
      store float %load42, ptr %faber.array.value41
      %faber.array.status41 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value41)
      %faber.old.status41 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error41 = icmp ne i32 %faber.old.status41, 0
      %faber.latched.status41 = select i1 %faber.has.error41, i32 %faber.old.status41, i32 %faber.array.status41
      store i32 %faber.latched.status41, ptr @__faber_rt_v1_status
      %load44 = load float, ptr %t12.addr
      %faber.array.value43 = alloca float
      store float %load44, ptr %faber.array.value43
      %faber.array.status43 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value43)
      %faber.old.status43 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error43 = icmp ne i32 %faber.old.status43, 0
      %faber.latched.status43 = select i1 %faber.has.error43, i32 %faber.old.status43, i32 %faber.array.status43
      store i32 %faber.latched.status43, ptr @__faber_rt_v1_status
      %load46 = load float, ptr %t13.addr
      %faber.array.value45 = alloca float
      store float %load46, ptr %faber.array.value45
      %faber.array.status45 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value45)
      %faber.old.status45 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error45 = icmp ne i32 %faber.old.status45, 0
      %faber.latched.status45 = select i1 %faber.has.error45, i32 %faber.old.status45, i32 %faber.array.status45
      store i32 %faber.latched.status45, ptr @__faber_rt_v1_status
      %load48 = load float, ptr %t14.addr
      %faber.array.value47 = alloca float
      store float %load48, ptr %faber.array.value47
      %faber.array.status47 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value47)
      %faber.old.status47 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error47 = icmp ne i32 %faber.old.status47, 0
      %faber.latched.status47 = select i1 %faber.has.error47, i32 %faber.old.status47, i32 %faber.array.status47
      store i32 %faber.latched.status47, ptr @__faber_rt_v1_status
      %load50 = load float, ptr %t15.addr
      %faber.array.value49 = alloca float
      store float %load50, ptr %faber.array.value49
      %faber.array.status49 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value49)
      %faber.old.status49 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error49 = icmp ne i32 %faber.old.status49, 0
      %faber.latched.status49 = select i1 %faber.has.error49, i32 %faber.old.status49, i32 %faber.array.status49
      store i32 %faber.latched.status49, ptr @__faber_rt_v1_status
      %load52 = load float, ptr %t16.addr
      %faber.array.value51 = alloca float
      store float %load52, ptr %faber.array.value51
      %faber.array.status51 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value51)
      %faber.old.status51 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error51 = icmp ne i32 %faber.old.status51, 0
      %faber.latched.status51 = select i1 %faber.has.error51, i32 %faber.old.status51, i32 %faber.array.status51
      store i32 %faber.latched.status51, ptr @__faber_rt_v1_status
      %load54 = load float, ptr %t17.addr
      %faber.array.value53 = alloca float
      store float %load54, ptr %faber.array.value53
      %faber.array.status53 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value53)
      %faber.old.status53 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error53 = icmp ne i32 %faber.old.status53, 0
      %faber.latched.status53 = select i1 %faber.has.error53, i32 %faber.old.status53, i32 %faber.array.status53
      store i32 %faber.latched.status53, ptr @__faber_rt_v1_status
      %load56 = load float, ptr %t18.addr
      %faber.array.value55 = alloca float
      store float %load56, ptr %faber.array.value55
      %faber.array.status55 = call i32 @__faber_rt_v1_array_push(ptr %faber.context32, ptr %faber.array.handle32, i32 5, ptr %faber.array.value55)
      %faber.old.status55 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error55 = icmp ne i32 %faber.old.status55, 0
      %faber.latched.status55 = select i1 %faber.has.error55, i32 %faber.old.status55, i32 %faber.array.status55
      store i32 %faber.latched.status55, ptr @__faber_rt_v1_status
      store ptr %faber.array.handle32, ptr %t19.addr
      %load57 = load ptr, ptr %t19.addr
      store ptr %load57, ptr %l1.addr
      %faber.context58 = load ptr, ptr @__faber_rt_v1_context
      %faber.tensor.result58 = call %FaberRtPtrResultV1 @__faber_rt_v1_tensor_new(ptr %faber.context58, i32 5)
      %faber.tensor.status58 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result58, 0
      %faber.tensor.value58 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result58, 1
      %faber.old.status58 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error58 = icmp ne i32 %faber.old.status58, 0
      %faber.latched.status58 = select i1 %faber.has.error58, i32 %faber.old.status58, i32 %faber.tensor.status58
      store i32 %faber.latched.status58, ptr @__faber_rt_v1_status
      store ptr %faber.tensor.value58, ptr %t20.addr
      %load59 = load ptr, ptr %t20.addr
      store ptr %load59, ptr %l2.addr
      %faber.context60 = load ptr, ptr @__faber_rt_v1_context
      %faber.array.result60 = call %FaberRtPtrResultV1 @__faber_rt_v1_array_new(ptr %faber.context60, i32 4)
      %faber.array.status60 = extractvalue %FaberRtPtrResultV1 %faber.array.result60, 0
      %faber.array.handle60 = extractvalue %FaberRtPtrResultV1 %faber.array.result60, 1
      %faber.old.status60 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error60 = icmp ne i32 %faber.old.status60, 0
      %faber.latched.status60 = select i1 %faber.has.error60, i32 %faber.old.status60, i32 %faber.array.status60
      store i32 %faber.latched.status60, ptr @__faber_rt_v1_status
      %faber.array.value61 = alloca i64
      store i64 2, ptr %faber.array.value61
      %faber.array.status61 = call i32 @__faber_rt_v1_array_push(ptr %faber.context60, ptr %faber.array.handle60, i32 4, ptr %faber.array.value61)
      %faber.old.status61 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error61 = icmp ne i32 %faber.old.status61, 0
      %faber.latched.status61 = select i1 %faber.has.error61, i32 %faber.old.status61, i32 %faber.array.status61
      store i32 %faber.latched.status61, ptr @__faber_rt_v1_status
      %faber.array.value62 = alloca i64
      store i64 3, ptr %faber.array.value62
      %faber.array.status62 = call i32 @__faber_rt_v1_array_push(ptr %faber.context60, ptr %faber.array.handle60, i32 4, ptr %faber.array.value62)
      %faber.old.status62 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error62 = icmp ne i32 %faber.old.status62, 0
      %faber.latched.status62 = select i1 %faber.has.error62, i32 %faber.old.status62, i32 %faber.array.status62
      store i32 %faber.latched.status62, ptr @__faber_rt_v1_status
      store ptr %faber.array.handle60, ptr %t21.addr
      %load64 = load ptr, ptr %l2.addr
      %faber.context63 = load ptr, ptr @__faber_rt_v1_context
      %load65 = load ptr, ptr %l0.addr
      %load66 = load ptr, ptr %t21.addr
      %faber.tensor.result63 = call %FaberRtPtrResultV1 @__faber_rt_v1_tensor_from_flat(ptr %faber.context63, i32 5, ptr %load65, ptr %load66)
      %faber.tensor.status63 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result63, 0
      %faber.tensor.value63 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result63, 1
      %faber.old.status63 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error63 = icmp ne i32 %faber.old.status63, 0
      %faber.latched.status63 = select i1 %faber.has.error63, i32 %faber.old.status63, i32 %faber.tensor.status63
      store i32 %faber.latched.status63, ptr @__faber_rt_v1_status
      store ptr %faber.tensor.value63, ptr %t22.addr
      %load67 = load ptr, ptr %t22.addr
      store ptr %load67, ptr %l3.addr
      %faber.context68 = load ptr, ptr @__faber_rt_v1_context
      %faber.array.result68 = call %FaberRtPtrResultV1 @__faber_rt_v1_array_new(ptr %faber.context68, i32 4)
      %faber.array.status68 = extractvalue %FaberRtPtrResultV1 %faber.array.result68, 0
      %faber.array.handle68 = extractvalue %FaberRtPtrResultV1 %faber.array.result68, 1
      %faber.old.status68 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error68 = icmp ne i32 %faber.old.status68, 0
      %faber.latched.status68 = select i1 %faber.has.error68, i32 %faber.old.status68, i32 %faber.array.status68
      store i32 %faber.latched.status68, ptr @__faber_rt_v1_status
      %faber.array.value69 = alloca i64
      store i64 3, ptr %faber.array.value69
      %faber.array.status69 = call i32 @__faber_rt_v1_array_push(ptr %faber.context68, ptr %faber.array.handle68, i32 4, ptr %faber.array.value69)
      %faber.old.status69 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error69 = icmp ne i32 %faber.old.status69, 0
      %faber.latched.status69 = select i1 %faber.has.error69, i32 %faber.old.status69, i32 %faber.array.status69
      store i32 %faber.latched.status69, ptr @__faber_rt_v1_status
      %faber.array.value70 = alloca i64
      store i64 4, ptr %faber.array.value70
      %faber.array.status70 = call i32 @__faber_rt_v1_array_push(ptr %faber.context68, ptr %faber.array.handle68, i32 4, ptr %faber.array.value70)
      %faber.old.status70 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error70 = icmp ne i32 %faber.old.status70, 0
      %faber.latched.status70 = select i1 %faber.has.error70, i32 %faber.old.status70, i32 %faber.array.status70
      store i32 %faber.latched.status70, ptr @__faber_rt_v1_status
      store ptr %faber.array.handle68, ptr %t23.addr
      %load72 = load ptr, ptr %l2.addr
      %faber.context71 = load ptr, ptr @__faber_rt_v1_context
      %load73 = load ptr, ptr %l1.addr
      %load74 = load ptr, ptr %t23.addr
      %faber.tensor.result71 = call %FaberRtPtrResultV1 @__faber_rt_v1_tensor_from_flat(ptr %faber.context71, i32 5, ptr %load73, ptr %load74)
      %faber.tensor.status71 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result71, 0
      %faber.tensor.value71 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result71, 1
      %faber.old.status71 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error71 = icmp ne i32 %faber.old.status71, 0
      %faber.latched.status71 = select i1 %faber.has.error71, i32 %faber.old.status71, i32 %faber.tensor.status71
      store i32 %faber.latched.status71, ptr @__faber_rt_v1_status
      store ptr %faber.tensor.value71, ptr %t24.addr
      %load75 = load ptr, ptr %t24.addr
      store ptr %load75, ptr %l4.addr
      %load77 = load ptr, ptr %l3.addr
      %faber.context76 = load ptr, ptr @__faber_rt_v1_context
      %load78 = load ptr, ptr %l4.addr
      %faber.tensor.result76 = call %FaberRtPtrResultV1 @__faber_rt_v1_tensor_matmul(ptr %faber.context76, ptr %load77, ptr %load78)
      %faber.tensor.status76 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result76, 0
      %faber.tensor.value76 = extractvalue %FaberRtPtrResultV1 %faber.tensor.result76, 1
      %faber.old.status76 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error76 = icmp ne i32 %faber.old.status76, 0
      %faber.latched.status76 = select i1 %faber.has.error76, i32 %faber.old.status76, i32 %faber.tensor.status76
      store i32 %faber.latched.status76, ptr @__faber_rt_v1_status
      store ptr %faber.tensor.value76, ptr %t25.addr
      %load79 = load ptr, ptr %t25.addr
      store ptr %load79, ptr %l5.addr
      %load81 = load ptr, ptr %l5.addr
      %faber.context80 = load ptr, ptr @__faber_rt_v1_context
      %faber.tensor.reduce.out80 = alloca float
      store float 0.0, ptr %faber.tensor.reduce.out80
      %faber.tensor.status80 = call i32 @__faber_rt_v1_tensor_mean(ptr %faber.context80, ptr %load81, i32 5, ptr %faber.tensor.reduce.out80)
      %faber.tensor.reduce80 = load float, ptr %faber.tensor.reduce.out80
      %faber.old.status80 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error80 = icmp ne i32 %faber.old.status80, 0
      %faber.latched.status80 = select i1 %faber.has.error80, i32 %faber.old.status80, i32 %faber.tensor.status80
      store i32 %faber.latched.status80, ptr @__faber_rt_v1_status
      store float %faber.tensor.reduce80, ptr %t26.addr
      %load82 = load float, ptr %t26.addr
      store float %load82, ptr %l6.addr
      %load83 = load float, ptr %l6.addr
      %faber.context84 = load ptr, ptr @__faber_rt_v1_context
      %faber.diag.status84 = call i32 @__faber_rt_v1_diagnostic_nota_f32(ptr %faber.context84, float %load83)
      %faber.old.status84 = load i32, ptr @__faber_rt_v1_status
      %faber.has.error84 = icmp ne i32 %faber.old.status84, 0
      %faber.latched.status84 = select i1 %faber.has.error84, i32 %faber.old.status84, i32 %faber.diag.status84
      store i32 %faber.latched.status84, ptr @__faber_rt_v1_status
      ret void
}

define %FaberRtExitV1 @__faber_program_entry_v1(ptr %context) {
    entry:
      store ptr %context, ptr @__faber_rt_v1_context
      call void @incipit()
      %faber.entry.status = load i32, ptr @__faber_rt_v1_status
      %faber.entry.status.ext = zext i32 %faber.entry.status to i64
      %faber.entry.shifted = shl i64 %faber.entry.status.ext, 32
      %faber.entry.packed = or i64 0, %faber.entry.shifted
      ret %FaberRtExitV1 %faber.entry.packed
}

// Generated by radix metal-text (supported-with-limitations compute source).
#include <metal_stdlib>
using namespace metal;

kernel void multiplico(
    device const float* a_in [[buffer(0)]],
    device const float* b_in [[buffer(1)]],
    device float* output [[buffer(2)]],
    uint3 id [[thread_position_in_grid]],
    uint3 local_id [[thread_position_in_threadgroup]]
) {
    uint i = id.x;
threadgroup float shared_a[64];
threadgroup float shared_b[64];
float acc = 0.0;
uint row = id.y;
uint col = id.x;
uint ty = local_id.y;
uint tx = local_id.x;
for (uint k_tile = 0u; k_tile < 1u; k_tile++) {
    uint k_start = k_tile * 8u;
    uint a_idx = row * 8u + (k_start + tx);
    if (a_idx < 16u * 8u) { shared_a[ty * 8u + tx] = a_in[a_idx]; }
    if (a_idx >= 16u * 8u) { shared_a[ty * 8u + tx] = 0.0; }
    uint b_idx = (k_start + ty) * 16u + col;
    if (b_idx < 8u * 16u) { shared_b[ty * 8u + tx] = b_in[b_idx]; }
    if (b_idx >= 8u * 16u) { shared_b[ty * 8u + tx] = 0.0; }
    threadgroup_barrier(mem_flags::mem_threadgroup);
    for (uint kk = 0u; kk < 8u; kk++) {
        acc += shared_a[ty * 8u + kk] * shared_b[kk * 8u + tx];
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);
}
uint out_idx = row * 16u + col;
if (row < 16u && col < 16u) { output[out_idx] = acc; }
}