[BUG] GPU compute pipeline creation fails: unlowered intrinsic `agc.simdgroup_matrix_8x8_load_bounds_checked_none_sz0_device_sz64`

Bug description

Actual behavior

When using external_call to invoke air.simdgroup_matrix_8x8_load / air.simdgroup_matrix_8x8_store / air.simdgroup_matrix_8x8_multiply_accumulate intrinsics inside a Mojo GPU kernel targeting Metal M4, compute pipeline creation fails at runtime with the following error:

Failed to create compute pipeline state:
Encountered unlowered function call to
agc.simdgroup_matrix_8x8_load_bounds_checked_none_sz0_device_sz64

This happens even though the generated AIR/LLVM IR is valid, and the kernel contains only calls to the official AIR intrinsics.

It appears that the AGC → AIR lowering pass is incomplete, Mojo internally emits an AGC intrinsic named agc.simdgroup_matrix_8x8_load_bounds_checked_none_sz0_device_sz64, which is never lowered, causing Metal pipeline creation to fail.

Expected behavior

The kernel should compile and run as long as the developer explicitly calls only:

• air.simdgroup_matrix_8x8_load
• air.simdgroup_matrix_8x8_store
• air.simdgroup_matrix_8x8_multiply_accumulate

The internal AGC intrinsic should be lowered to AIR properly, or should not be emitted.

Steps to reproduce

from sys import external_call


fn _get_simdgroup_suffix[dtype: DType]() -> String:
    constrained[
        dtype in (DType.float16, DType.float32),
        "unsupported dtype for simdgroup operations",
    ]()
    return "f16" if dtype == DType.float16 else "f32"


fn simdgroup_load[
    dtype: DType, //, transpose: Bool = False
](ptr: UnsafePointer[Scalar[dtype], **_], stride: Int) -> SIMD[dtype, 64]:
    comptime dtype_suffix = _get_simdgroup_suffix[dtype]()
    comptime resolved_addr_space = AddressSpace.GLOBAL if ptr.address_space == AddressSpace.GENERIC else ptr.address_space
    comptime addr_space_code = "1" if resolved_addr_space == AddressSpace.GLOBAL else "3"
    comptime intrin = "air.simdgroup_matrix_8x8_load.v64" + dtype_suffix + ".p" + addr_space_code + dtype_suffix
    return external_call[intrin, SIMD[dtype, 64]](
        ptr.address_space_cast[resolved_addr_space](),
        stride,
        SIMD[DType.int64, 2](0, 0),
        transpose,
    )


fn simdgroup_store[
    dtype: DType, //, transpose: Bool = False
](ptr: UnsafePointer[Scalar[dtype], **_], value: SIMD[dtype, 64], stride: Int):
    comptime dtype_suffix = _get_simdgroup_suffix[dtype]()
    comptime resolved_addr_space = AddressSpace.GLOBAL if ptr.address_space == AddressSpace.GENERIC else ptr.address_space
    comptime addr_space_code = "1" if resolved_addr_space == AddressSpace.GLOBAL else "3"
    comptime intrin = "air.simdgroup_matrix_8x8_store.v64" + dtype_suffix + ".p" + addr_space_code + dtype_suffix
    external_call[intrin, NoneType](
        value,
        ptr.address_space_cast[resolved_addr_space](),
        stride,
        SIMD[DType.int64, 2](0, 0),
        transpose,
    )


fn simdgroup_multiply_accumulate[dtype: DType](
    acc: SIMD[dtype, 64],
    a: SIMD[dtype, 64],
    b: SIMD[dtype, 64],
) -> SIMD[dtype, 64]:
    comptime dtype_suffix = _get_simdgroup_suffix[dtype]() 
    comptime intrin = "air.simdgroup_matrix_8x8_multiply_accumulate" + (".v64" + dtype_suffix) * 4
    return external_call[intrin, SIMD[dtype, 64]](a, b, acc)


from gpu.host.info import MetalM4
from gpu.host.compile import _compile_code
from gpu.host import DeviceContext
from layout import LayoutTensor, Layout

fn kernel[dtype: DType](
    out_ptr: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a_ptr: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    b_ptr: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    c_ptr: UnsafePointer[Scalar[dtype], MutAnyOrigin],
):
    var a = simdgroup_load(a_ptr, 8)
    var b = simdgroup_load(b_ptr, 8)
    var c = simdgroup_load(c_ptr, 8)

   # var r = simdgroup_multiply_accumulate(c, a, b)
    simdgroup_store(out_ptr, c, 8)

def main():
    print(
        _compile_code[
            kernel[DType.float32],
            target = MetalM4.target(),
            emission_kind="llvm-opt",
        ]()
    )
    var ctx = DeviceContext()
    var a = ctx.enqueue_create_buffer[DType.float32](64)
    var b = ctx.enqueue_create_buffer[DType.float32](64)
    var c = ctx.enqueue_create_buffer[DType.float32](64)
    var out = ctx.enqueue_create_buffer[DType.float32](64)
    a.enqueue_fill(1.0)
    b.enqueue_fill(2.0)
    c.enqueue_fill(3.0)
    ctx.enqueue_function_checked[
        kernel[DType.float32],
        kernel[DType.float32],
    ](
        out,
        a,
        b,
        c,
        grid_dim=1,
        block_dim=32,
    )
    ctx.synchronize()
    with out.map_to_host() as host_out:
        for i in range(64):
            print(host_out[i])
    

Unhandled exception caught during execution:
Failed to create compute pipeline state:
Encountered unlowered function call to agc.simdgroup_matrix_8x8_load_bounds_checked_none_sz0_device_sz64

System information

System
------------
       Pixi version: 0.56.0
           Platform: osx-arm64
   Virtual packages: __unix=0=0
                   : __osx=26.0.1=0
                   : __archspec=1=m1

Environments
------------
        Environment: default
           Features: default
           Channels: https://conda.modular.com/max-nightly, conda-forge
   Dependency count: 2
       Dependencies: mojo, max
   Target platforms: osx-arm64

Package   Version               Build        Size      Kind   Source
max       26.1.0.dev2025111405  3.14release  9.2 MiB   conda  https://conda.modular.com/max-nightly/
max-core  26.1.0.dev2025111405  release      72.2 MiB  conda  https://conda.modular.com/max-nightly/

Reproduce the failure with the provided Mojo kernel, _compile_code, and MetalM4.target(), then inspect the generated AIR/LLVM IR and the AGC-to-AIR lowering path. Done means the internal agc.simdgroup_matrix_8x8_load_bounds_checked_none_sz0_device_sz64 call is lowered or omitted and the Metal compute pipeline is created successfully.