`sincos` intrinsic fails to compile with Float16

[christiangnrd]

MWE:

using Metal
mtl_a = Metal.ones(Float16, 1)
mtl_b = copy(mtl_a)
function intr_test3(arr_sin, arr_cos)
    idx = thread_position_in_grid_1d()
    s, c = sincos(arr_cos[idx])
    arr_sin[idx] = s
    arr_cos[idx] = c
    return nothing
end
@metal intr_test3(mtl_a, mtl_b) # Fails here

Stacktrace

ERROR: Compilation to native code failed; see below for details.
If you think this is a bug, please file an issue and attach the following files:
- /var/folders/4g/lnkpkf3s4rxd_wbl8vwnqs4r0000gn/T/jl_Al0dKrkbiL.ll
- /var/folders/4g/lnkpkf3s4rxd_wbl8vwnqs4r0000gn/T/jl_wrUlMvonIP.air
- /var/folders/4g/lnkpkf3s4rxd_wbl8vwnqs4r0000gn/T/jl_mt69yHWJcV.metallib
Stacktrace:
  [1] error(s::String)
    @ Base ./error.jl:35
  [2] macro expansion
    @ ~/.julia/dev/Metal/src/compiler/compilation.jl:205 [inlined]
  [3] macro expansion
    @ ~/.julia/packages/ObjectiveC/CUn6a/src/os.jl:264 [inlined]
  [4] macro expansion
    @ ~/.julia/dev/Metal/src/compiler/compilation.jl:182 [inlined]
  [5] (::Metal.var"#173#174"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, @NamedTuple{ir::String, air::Vector{UInt8}, metallib::Vector{UInt8}, entry::String}})()
    @ Metal ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:644
  [6] macro expansion
    @ ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:572 [inlined]
  [7] macro expansion
    @ ./lock.jl:273 [inlined]
  [8] ObjectiveC.Foundation.NSAutoreleasePool(f::Metal.var"#173#174"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, @NamedTuple{ir::String, air::Vector{UInt8}, metallib::Vector{UInt8}, entry::String}})
    @ ObjectiveC.Foundation ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:564
  [9] link(job::GPUCompiler.CompilerJob, compiled::@NamedTuple{ir::String, air::Vector{UInt8}, metallib::Vector{UInt8}, entry::String})
    @ Metal ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:643
 [10] actual_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, world::UInt64, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/.julia/packages/GPUCompiler/Nxf8r/src/execution.jl:262
 [11] cached_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::Function, linker::Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/Nxf8r/src/execution.jl:151
 [12] macro expansion
    @ ~/.julia/dev/Metal/src/compiler/execution.jl:189 [inlined]
 [13] macro expansion
    @ ./lock.jl:273 [inlined]
 [14] mtlfunction(f::typeof(intr_test3), tt::Type{Tuple{MtlDeviceVector{Float16, 1}, MtlDeviceVector{Float16, 1}}}; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/.julia/dev/Metal/src/compiler/execution.jl:184
 [15] mtlfunction(f::typeof(intr_test3), tt::Type{Tuple{MtlDeviceVector{Float16, 1}, MtlDeviceVector{Float16, 1}}})
    @ Metal ~/.julia/dev/Metal/src/compiler/execution.jl:182
 [16] top-level scope
    @ ~/.julia/dev/Metal/src/compiler/execution.jl:85

caused by: NSError: Compiler encountered an internal error (AGXMetalG14X, code 3)
Stacktrace:
  [1] Metal.MTL.MTLComputePipelineState(dev::Metal.MTL.MTLDeviceInstance, fun::Metal.MTL.MTLFunctionInstance)
    @ Metal.MTL ~/.julia/dev/Metal/lib/mtl/compute_pipeline.jl:28
  [2] macro expansion
    @ ~/.julia/dev/Metal/src/compiler/compilation.jl:187 [inlined]
  [3] macro expansion
    @ ~/.julia/packages/ObjectiveC/CUn6a/src/os.jl:264 [inlined]
  [4] macro expansion
    @ ~/.julia/dev/Metal/src/compiler/compilation.jl:182 [inlined]
  [5] (::Metal.var"#173#174"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, @NamedTuple{ir::String, air::Vector{UInt8}, metallib::Vector{UInt8}, entry::String}})()
    @ Metal ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:644
  [6] macro expansion
    @ ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:572 [inlined]
  [7] macro expansion
    @ ./lock.jl:273 [inlined]
  [8] ObjectiveC.Foundation.NSAutoreleasePool(f::Metal.var"#173#174"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, @NamedTuple{ir::String, air::Vector{UInt8}, metallib::Vector{UInt8}, entry::String}})
    @ ObjectiveC.Foundation ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:564
  [9] link(job::GPUCompiler.CompilerJob, compiled::@NamedTuple{ir::String, air::Vector{UInt8}, metallib::Vector{UInt8}, entry::String})
    @ Metal ~/.julia/packages/ObjectiveC/CUn6a/src/foundation.jl:643
 [10] actual_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, world::UInt64, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/.julia/packages/GPUCompiler/Nxf8r/src/execution.jl:262
 [11] cached_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::Function, linker::Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/Nxf8r/src/execution.jl:151
 [12] macro expansion
    @ ~/.julia/dev/Metal/src/compiler/execution.jl:189 [inlined]
 [13] macro expansion
    @ ./lock.jl:273 [inlined]
 [14] mtlfunction(f::typeof(intr_test3), tt::Type{Tuple{MtlDeviceVector{Float16, 1}, MtlDeviceVector{Float16, 1}}}; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/.julia/dev/Metal/src/compiler/execution.jl:184
 [15] mtlfunction(f::typeof(intr_test3), tt::Type{Tuple{MtlDeviceVector{Float16, 1}, MtlDeviceVector{Float16, 1}}})
    @ Metal ~/.julia/dev/Metal/src/compiler/execution.jl:182
 [16] top-level scope
    @ ~/.julia/dev/Metal/src/compiler/execution.jl:85

IR file

; ModuleID = 'start'
source_filename = "start"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-n8:16:32"
target triple = "air64-apple-macosx15.3.0"

; Function Attrs: alwaysinline cold nounwind
declare void @llvm.trap() #0

declare half @air.sincos.f16(half, i64) local_unnamed_addr

define internal fastcc void @gpu_report_exception() unnamed_addr !dbg !50 {
top:
  ret void, !dbg !53
}

define internal fastcc void @gpu_signal_exception() unnamed_addr !dbg !54 {
top:
  ret void, !dbg !56
}

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1

define void @_Z10intr_test314MtlDeviceArrayI7Float16Li1ELi1EES1_({ i8 addrspace(1)*, [1 x i64] } addrspace(1)* %0, { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %1, i32 %thread_position_in_grid) local_unnamed_addr !dbg !57 {
conversion:
  %2 = bitcast { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %0 to half addrspace(1)* addrspace(1)*
  %.unpack10 = load half addrspace(1)*, half addrspace(1)* addrspace(1)* %2, align 8
  %3 = getelementptr inbounds { i8 addrspace(1)*, [1 x i64] }, { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %1, i64 0, i32 1, i64 0
  %.unpack12.unpack = load i64, i64 addrspace(1)* %3, align 8
  %4 = alloca i16, align 4
  %5 = add i32 %thread_position_in_grid, 1, !dbg !59
  %6 = icmp ne i32 %5, 0, !dbg !66
  %7 = icmp sgt i64 %.unpack12.unpack, -1, !dbg !84
  %8 = zext i32 %5 to i64, !dbg !89
  %9 = icmp uge i64 %.unpack12.unpack, %8, !dbg !102
  %10 = and i1 %7, %9, !dbg !103
  %11 = and i1 %6, %10, !dbg !106
  br i1 %11, label %L38, label %L20, !dbg !77

L20:                                              ; preds = %conversion
  call fastcc void @gpu_report_exception(), !dbg !107
  call fastcc void @gpu_signal_exception(), !dbg !107
  call void @llvm.trap(), !dbg !107
  unreachable, !dbg !107

L38:                                              ; preds = %conversion
  %12 = bitcast { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %1 to half addrspace(1)* addrspace(1)*
  %.unpack14 = load half addrspace(1)*, half addrspace(1)* addrspace(1)* %12, align 8
  %13 = getelementptr inbounds { i8 addrspace(1)*, [1 x i64] }, { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %0, i64 0, i32 1, i64 0
  %.unpack8.unpack = load i64, i64 addrspace(1)* %13, align 8
  %14 = sext i32 %thread_position_in_grid to i64, !dbg !111
  %15 = getelementptr inbounds half, half addrspace(1)* %.unpack14, i64 %14, !dbg !111
  %16 = load half, half addrspace(1)* %15, align 2, !dbg !111, !tbaa !123
  store i16 0, i16* %4, align 4, !dbg !126
  %17 = ptrtoint i16* %4 to i64, !dbg !135
  %18 = call half @air.sincos.f16(half %16, i64 %17), !dbg !140
  %19 = icmp slt i64 %.unpack8.unpack, 0, !dbg !141
  %20 = icmp ult i64 %.unpack8.unpack, %8, !dbg !152
  %.not2 = or i1 %19, %20, !dbg !146
  br i1 %.not2, label %L60, label %L89, !dbg !146

L60:                                              ; preds = %L38
  call fastcc void @gpu_report_exception(), !dbg !154
  call fastcc void @gpu_signal_exception(), !dbg !154
  call void @llvm.trap(), !dbg !154
  unreachable, !dbg !154

L89:                                              ; preds = %L38
  %21 = bitcast i16* %4 to half*, !dbg !156
  %22 = load half, half* %21, align 4, !dbg !156, !tbaa !162, !alias.scope !168, !noalias !171
  %23 = getelementptr inbounds half, half addrspace(1)* %.unpack10, i64 %14, !dbg !176
  store half %18, half addrspace(1)* %23, align 2, !dbg !176, !tbaa !123
  store half %22, half addrspace(1)* %15, align 2, !dbg !183, !tbaa !123
  ret void, !dbg !190
}

attributes #0 = { alwaysinline cold nounwind }
attributes #1 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }

!llvm.module.flags = !{!0, !1, !2, !3, !4, !5, !6, !7, !8}
!llvm.dbg.cu = !{!9, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39}
!julia.kernel = !{!40}
!air.kernel = !{!41}
!llvm.ident = !{!47}
!air.version = !{!48}
!air.language_version = !{!49}

!0 = !{i32 2, !"Dwarf Version", i32 4}
!1 = !{i32 2, !"Debug Info Version", i32 3}
!2 = !{i32 7, !"air.max_device_buffers", i32 31}
!3 = !{i32 7, !"air.max_constant_buffers", i32 31}
!4 = !{i32 7, !"air.max_threadgroup_buffers", i32 31}
!5 = !{i32 7, !"air.max_textures", i32 128}
!6 = !{i32 7, !"air.max_read_write_textures", i32 8}
!7 = !{i32 7, !"air.max_samplers", i32 16}
!8 = !{i32 2, !"SDK Version", [2 x i32] [i32 15, i32 3]}
!9 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!10 = !DIFile(filename: "julia", directory: ".")
!11 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!12 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!13 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!14 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!15 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!16 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!17 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!18 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!19 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!20 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!21 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!22 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!23 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!24 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!25 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!26 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!27 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!28 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!29 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!30 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!31 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!32 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!33 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!34 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!35 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!36 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!37 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!38 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!39 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly, nameTableKind: None)
!40 = !{void ({ i8 addrspace(1)*, [1 x i64] } addrspace(1)*, { i8 addrspace(1)*, [1 x i64] } addrspace(1)*, i32)* @_Z10intr_test314MtlDeviceArrayI7Float16Li1ELi1EES1_}
!41 = !{void ({ i8 addrspace(1)*, [1 x i64] } addrspace(1)*, { i8 addrspace(1)*, [1 x i64] } addrspace(1)*, i32)* @_Z10intr_test314MtlDeviceArrayI7Float16Li1ELi1EES1_, !42, !43}
!42 = !{}
!43 = !{!44, !45, !46}
!44 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 16, !"air.arg_type_align_size", i32 8, !"air.arg_type_name", !"MtlDeviceVector{Float16, 1}", !"air.arg_name", !"arr_sin"}
!45 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 16, !"air.arg_type_align_size", i32 8, !"air.arg_type_name", !"MtlDeviceVector{Float16, 1}", !"air.arg_name", !"arr_cos"}
!46 = !{i32 2, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint"}
!47 = !{!"Julia 1.11.3 with Metal.jl"}
!48 = !{i32 2, i32 5, i32 0}
!49 = !{!"Metal", i32 3, i32 2, i32 0}
!50 = distinct !DISubprogram(name: "report_exception", linkageName: "julia_report_exception_12485", scope: null, file: !51, line: 13, type: !52, scopeLine: 13, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !15, retainedNodes: !42)
!51 = !DIFile(filename: "/Users/christian/.julia/dev/Metal/src/device/runtime.jl", directory: ".")
!52 = !DISubroutineType(cc: DW_CC_nocall, types: !42)
!53 = !DILocation(line: 13, scope: !50)
!54 = distinct !DISubprogram(name: "signal_exception", linkageName: "julia_signal_exception_12557", scope: null, file: !51, line: 9, type: !55, scopeLine: 9, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !23, retainedNodes: !42)
!55 = !DISubroutineType(types: !42)
!56 = !DILocation(line: 9, scope: !54)
!57 = distinct !DISubprogram(name: "intr_test3", linkageName: "julia_intr_test3_19106", scope: null, file: !58, line: 1, type: !55, scopeLine: 1, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!58 = !DIFile(filename: "REPL[1]", directory: ".")
!59 = !DILocation(line: 87, scope: !60, inlinedAt: !62)
!60 = distinct !DISubprogram(name: "+;", linkageName: "+", scope: !61, file: !61, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!61 = !DIFile(filename: "int.jl", directory: ".")
!62 = !DILocation(line: 49, scope: !63, inlinedAt: !65)
!63 = distinct !DISubprogram(name: "#thread_position_in_grid_1d;", linkageName: "#thread_position_in_grid_1d", scope: !64, file: !64, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!64 = !DIFile(filename: "/Users/christian/.julia/dev/Metal/src/device/intrinsics/arguments.jl", directory: ".")
!65 = !DILocation(line: 2, scope: !57)
!66 = !DILocation(line: 515, scope: !67, inlinedAt: !68)
!67 = distinct !DISubprogram(name: "<=;", linkageName: "<=", scope: !61, file: !61, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!68 = !DILocation(line: 485, scope: !69, inlinedAt: !71)
!69 = distinct !DISubprogram(name: "<=;", linkageName: "<=", scope: !70, file: !70, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!70 = !DIFile(filename: "promotion.jl", directory: ".")
!71 = !DILocation(line: 521, scope: !67, inlinedAt: !72)
!72 = !DILocation(line: 752, scope: !73, inlinedAt: !75)
!73 = distinct !DISubprogram(name: "checkindex;", linkageName: "checkindex", scope: !74, file: !74, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!74 = !DIFile(filename: "abstractarray.jl", directory: ".")
!75 = !DILocation(line: 689, scope: !76, inlinedAt: !77)
!76 = distinct !DISubprogram(name: "checkbounds;", linkageName: "checkbounds", scope: !74, file: !74, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!77 = !DILocation(line: 699, scope: !76, inlinedAt: !78)
!78 = !DILocation(line: 80, scope: !79, inlinedAt: !81)
!79 = distinct !DISubprogram(name: "arrayref;", linkageName: "arrayref", scope: !80, file: !80, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!80 = !DIFile(filename: "/Users/christian/.julia/dev/Metal/src/device/array.jl", directory: ".")
!81 = !DILocation(line: 103, scope: !82, inlinedAt: !83)
!82 = distinct !DISubprogram(name: "getindex;", linkageName: "getindex", scope: !80, file: !80, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!83 = !DILocation(line: 3, scope: !57)
!84 = !DILocation(line: 514, scope: !67, inlinedAt: !85)
!85 = !DILocation(line: 426, scope: !86, inlinedAt: !88)
!86 = distinct !DISubprogram(name: ">=;", linkageName: ">=", scope: !87, file: !87, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!87 = !DIFile(filename: "operators.jl", directory: ".")
!88 = !DILocation(line: 522, scope: !67, inlinedAt: !72)
!89 = !DILocation(line: 871, scope: !90, inlinedAt: !92)
!90 = distinct !DISubprogram(name: "toUInt64;", linkageName: "toUInt64", scope: !91, file: !91, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!91 = !DIFile(filename: "boot.jl", directory: ".")
!92 = !DILocation(line: 897, scope: !93, inlinedAt: !94)
!93 = distinct !DISubprogram(name: "UInt64;", linkageName: "UInt64", scope: !91, file: !91, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!94 = !DILocation(line: 7, scope: !95, inlinedAt: !97)
!95 = distinct !DISubprogram(name: "convert;", linkageName: "convert", scope: !96, file: !96, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!96 = !DIFile(filename: "number.jl", directory: ".")
!97 = !DILocation(line: 375, scope: !98, inlinedAt: !99)
!98 = distinct !DISubprogram(name: "_promote;", linkageName: "_promote", scope: !70, file: !70, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!99 = !DILocation(line: 400, scope: !100, inlinedAt: !101)
!100 = distinct !DISubprogram(name: "promote;", linkageName: "promote", scope: !70, file: !70, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!101 = !DILocation(line: 485, scope: !69, inlinedAt: !88)
!102 = !DILocation(line: 515, scope: !67, inlinedAt: !101)
!103 = !DILocation(line: 38, scope: !104, inlinedAt: !88)
!104 = distinct !DISubprogram(name: "&;", linkageName: "&", scope: !105, file: !105, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!105 = !DIFile(filename: "bool.jl", directory: ".")
!106 = !DILocation(line: 38, scope: !104, inlinedAt: !72)
!107 = !DILocation(line: 4, scope: !108, inlinedAt: !110)
!108 = distinct !DISubprogram(name: "#throw_boundserror", linkageName: "julia_#throw_boundserror_19134", scope: null, file: !109, line: 44, type: !52, scopeLine: 44, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !11, retainedNodes: !42)
!109 = !DIFile(filename: "/Users/christian/.julia/dev/Metal/src/device/quirks.jl", directory: ".")
!110 = distinct !DILocation(line: 699, scope: !76, inlinedAt: !78)
!111 = !DILocation(line: 39, scope: !112, inlinedAt: !114)
!112 = distinct !DISubprogram(name: "macro expansion;", linkageName: "macro expansion", scope: !113, file: !113, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!113 = !DIFile(filename: "/Users/christian/.julia/packages/LLVM/b3kFs/src/interop/base.jl", directory: ".")
!114 = !DILocation(line: 0, scope: !115, inlinedAt: !117)
!115 = distinct !DISubprogram(name: "macro expansion;", linkageName: "macro expansion", scope: !116, file: !116, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!116 = !DIFile(filename: "none", directory: ".")
!117 = !DILocation(line: 0, scope: !118, inlinedAt: !119)
!118 = distinct !DISubprogram(name: "pointerref;", linkageName: "pointerref", scope: !116, file: !116, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!119 = !DILocation(line: 85, scope: !120, inlinedAt: !122)
!120 = distinct !DISubprogram(name: "unsafe_load;", linkageName: "unsafe_load", scope: !121, file: !121, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!121 = !DIFile(filename: "/Users/christian/.julia/packages/LLVM/b3kFs/src/interop/pointer.jl", directory: ".")
!122 = !DILocation(line: 82, scope: !79, inlinedAt: !81)
!123 = !{!124, !124, i64 0, i64 0}
!124 = !{!"custom_tbaa_addrspace(1)", !125, i64 0}
!125 = !{!"custom_tbaa"}
!126 = !DILocation(line: 7, scope: !127, inlinedAt: !129)
!127 = distinct !DISubprogram(name: "RefValue;", linkageName: "RefValue", scope: !128, file: !128, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!128 = !DIFile(filename: "refvalue.jl", directory: ".")
!129 = !DILocation(line: 144, scope: !130, inlinedAt: !132)
!130 = distinct !DISubprogram(name: "Ref;", linkageName: "Ref", scope: !131, file: !131, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!131 = !DIFile(filename: "refpointer.jl", directory: ".")
!132 = !DILocation(line: 244, scope: !133, inlinedAt: !83)
!133 = distinct !DISubprogram(name: "#sincos;", linkageName: "#sincos", scope: !134, file: !134, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!134 = !DIFile(filename: "/Users/christian/.julia/dev/Metal/src/device/intrinsics/math.jl", directory: ".")
!135 = !DILocation(line: 30, scope: !136, inlinedAt: !138)
!136 = distinct !DISubprogram(name: "convert;", linkageName: "convert", scope: !137, file: !137, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!137 = !DIFile(filename: "pointer.jl", directory: ".")
!138 = !DILocation(line: 53, scope: !139, inlinedAt: !140)
!139 = distinct !DISubprogram(name: "unsafe_convert;", linkageName: "unsafe_convert", scope: !128, file: !128, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!140 = !DILocation(line: 245, scope: !133, inlinedAt: !83)
!141 = !DILocation(line: 514, scope: !67, inlinedAt: !142)
!142 = !DILocation(line: 426, scope: !86, inlinedAt: !143)
!143 = !DILocation(line: 522, scope: !67, inlinedAt: !144)
!144 = !DILocation(line: 752, scope: !73, inlinedAt: !145)
!145 = !DILocation(line: 689, scope: !76, inlinedAt: !146)
!146 = !DILocation(line: 699, scope: !76, inlinedAt: !147)
!147 = !DILocation(line: 86, scope: !148, inlinedAt: !149)
!148 = distinct !DISubprogram(name: "arrayset;", linkageName: "arrayset", scope: !80, file: !80, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!149 = !DILocation(line: 105, scope: !150, inlinedAt: !151)
!150 = distinct !DISubprogram(name: "setindex!;", linkageName: "setindex!", scope: !80, file: !80, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!151 = !DILocation(line: 4, scope: !57)
!152 = !DILocation(line: 515, scope: !67, inlinedAt: !153)
!153 = !DILocation(line: 485, scope: !69, inlinedAt: !143)
!154 = !DILocation(line: 4, scope: !108, inlinedAt: !155)
!155 = distinct !DILocation(line: 699, scope: !76, inlinedAt: !147)
!156 = !DILocation(line: 49, scope: !157, inlinedAt: !159)
!157 = distinct !DISubprogram(name: "getproperty;", linkageName: "getproperty", scope: !158, file: !158, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!158 = !DIFile(filename: "Base.jl", directory: ".")
!159 = !DILocation(line: 59, scope: !160, inlinedAt: !161)
!160 = distinct !DISubprogram(name: "getindex;", linkageName: "getindex", scope: !128, file: !128, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!161 = !DILocation(line: 246, scope: !133, inlinedAt: !83)
!162 = !{!163, !163, i64 0}
!163 = !{!"jtbaa_mutab", !164, i64 0}
!164 = !{!"jtbaa_value", !165, i64 0}
!165 = !{!"jtbaa_data", !166, i64 0}
!166 = !{!"jtbaa", !167, i64 0}
!167 = !{!"jtbaa"}
!168 = !{!169}
!169 = !{!"jnoalias_data", !170}
!170 = !{!"jnoalias"}
!171 = !{!172, !173, !174, !175}
!172 = !{!"jnoalias_gcframe", !170}
!173 = !{!"jnoalias_stack", !170}
!174 = !{!"jnoalias_typemd", !170}
!175 = !{!"jnoalias_const", !170}
!176 = !DILocation(line: 39, scope: !112, inlinedAt: !177)
!177 = !DILocation(line: 0, scope: !115, inlinedAt: !178)
!178 = !DILocation(line: 0, scope: !179, inlinedAt: !180)
!179 = distinct !DISubprogram(name: "pointerset;", linkageName: "pointerset", scope: !116, file: !116, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!180 = !DILocation(line: 88, scope: !181, inlinedAt: !182)
!181 = distinct !DISubprogram(name: "unsafe_store!;", linkageName: "unsafe_store!", scope: !121, file: !121, type: !55, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !9, retainedNodes: !42)
!182 = !DILocation(line: 88, scope: !148, inlinedAt: !149)
!183 = !DILocation(line: 39, scope: !112, inlinedAt: !184)
!184 = !DILocation(line: 0, scope: !115, inlinedAt: !185)
!185 = !DILocation(line: 0, scope: !179, inlinedAt: !186)
!186 = !DILocation(line: 88, scope: !181, inlinedAt: !187)
!187 = !DILocation(line: 88, scope: !148, inlinedAt: !188)
!188 = !DILocation(line: 105, scope: !150, inlinedAt: !189)
!189 = !DILocation(line: 5, scope: !57)
!190 = !DILocation(line: 6, scope: !57)

[maleadt]

Should be easy enough to replace with something Julian, although it may be valuable to check first what Metal C emits for this (if even supported).

[christiangnrd]

The @air.sincos.f16 intrinsic seems to exist so the failure is very strange.

sincoshalf.metal

#include <metal_stdlib>

using namespace metal;

kernel void sincoshalf(device half* sin,
                          device half* cos,
                          device half* in,
                          uint i [[thread_position_in_grid]])
{
    half tmpcos;
    sin[i] = sincos(in[i], tmpcos);
    cos[i] = tmpcos;
}

$ xcrun metal -S -emit-llvm -fno-fast-math sincoshalf.metal

sincoshalf.ll

; ModuleID = 'sincoshalf.metal'
source_filename = "sincoshalf.metal"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-n8:16:32"
target triple = "air64-apple-macosx15.0.0"

; Function Attrs: argmemonly mustprogress nounwind willreturn
define void @sincoshalf(half addrspace(1)* nocapture noundef writeonly "air-buffer-no-alias" %0, half addrspace(1)* nocapture noundef writeonly "air-buffer-no-alias" %1, half addrspace(1)* nocapture noundef readonly "air-buffer-no-alias" %2, i32 noundef %3) local_unnamed_addr #0 {
  %5 = alloca half, align 2
  %6 = bitcast half* %5 to i8*
  call void @llvm.lifetime.start.p0i8(i64 2, i8* nonnull %6) #3
  %7 = zext i32 %3 to i64
  %8 = getelementptr inbounds half, half addrspace(1)* %2, i64 %7
  %9 = load half, half addrspace(1)* %8, align 2, !tbaa !23, !alias.scope !27, !noalias !30
  %10 = call half @air.sincos.f16(half %9, half* nocapture nonnull writeonly %5) #4
  %11 = getelementptr inbounds half, half addrspace(1)* %0, i64 %7
  store half %10, half addrspace(1)* %11, align 2, !tbaa !23, !alias.scope !33, !noalias !34
  %12 = load half, half* %5, align 2, !tbaa !23
  %13 = getelementptr inbounds half, half addrspace(1)* %1, i64 %7
  store half %12, half addrspace(1)* %13, align 2, !tbaa !23, !alias.scope !35, !noalias !36
  call void @llvm.lifetime.end.p0i8(i64 2, i8* nonnull %6) #3
  ret void
}

; Function Attrs: argmemonly mustprogress nocallback nofree nosync nounwind willreturn
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1

; Function Attrs: argmemonly mustprogress nocallback nofree nosync nounwind willreturn
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1

; Function Attrs: argmemonly mustprogress nounwind willreturn
declare half @air.sincos.f16(half, half* nocapture writeonly) local_unnamed_addr #2

attributes #0 = { argmemonly mustprogress nounwind willreturn "frame-pointer"="all" "min-legal-vector-width"="0" "no-builtins" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #1 = { argmemonly mustprogress nocallback nofree nosync nounwind willreturn }
attributes #2 = { argmemonly mustprogress nounwind willreturn }
attributes #3 = { nounwind }
attributes #4 = { argmemonly nounwind willreturn }

!llvm.module.flags = !{!0, !1, !2, !3, !4, !5, !6, !7, !8}
!air.kernel = !{!9}
!air.compile_options = !{!16, !17, !18}
!llvm.ident = !{!19}
!air.version = !{!20}
!air.language_version = !{!21}
!air.source_file_name = !{!22}

!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 15, i32 2]}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 7, !"frame-pointer", i32 2}
!3 = !{i32 7, !"air.max_device_buffers", i32 31}
!4 = !{i32 7, !"air.max_constant_buffers", i32 31}
!5 = !{i32 7, !"air.max_threadgroup_buffers", i32 31}
!6 = !{i32 7, !"air.max_textures", i32 128}
!7 = !{i32 7, !"air.max_read_write_textures", i32 8}
!8 = !{i32 7, !"air.max_samplers", i32 16}
!9 = !{void (half addrspace(1)*, half addrspace(1)*, half addrspace(1)*, i32)* @sincoshalf, !10, !11}
!10 = !{}
!11 = !{!12, !13, !14, !15}
!12 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 2, !"air.arg_type_align_size", i32 2, !"air.arg_type_name", !"half", !"air.arg_name", !"sin"}
!13 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 2, !"air.arg_type_align_size", i32 2, !"air.arg_type_name", !"half", !"air.arg_name", !"cos"}
!14 = !{i32 2, !"air.buffer", !"air.location_index", i32 2, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 2, !"air.arg_type_align_size", i32 2, !"air.arg_type_name", !"half", !"air.arg_name", !"in"}
!15 = !{i32 3, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint", !"air.arg_name", !"i"}
!16 = !{!"air.compile.denorms_disable"}
!17 = !{!"air.compile.fast_math_disable"}
!18 = !{!"air.compile.framebuffer_fetch_enable"}
!19 = !{!"Apple metal version 32023.404 (metalfe-32023.404)"}
!20 = !{i32 2, i32 7, i32 0}
!21 = !{!"Metal", i32 3, i32 2, i32 0}
!22 = !{!"/Users/christian/.julia/dev/Metal/sincoshalf.metal"}
!23 = !{!24, !24, i64 0}
!24 = !{!"half", !25, i64 0}
!25 = !{!"omnipotent char", !26, i64 0}
!26 = !{!"Simple C++ TBAA"}
!27 = !{!28}
!28 = distinct !{!28, !29, !"air-alias-scope-arg(2)"}
!29 = distinct !{!29, !"air-alias-scopes(sincoshalf)"}
!30 = !{!31, !32}
!31 = distinct !{!31, !29, !"air-alias-scope-arg(0)"}
!32 = distinct !{!32, !29, !"air-alias-scope-arg(1)"}
!33 = !{!31}
!34 = !{!32, !28}
!35 = !{!32}
!36 = !{!31, !28}

[maleadt]

  "termination": {
    "code": 1,
    "flags": 518,
    "namespace": "METAL",
    "reasons": [
      "unable to legalize instruction: 96 %56:_(s16), %33:_(s64) :: (store (s16) into %ir.22, align 4, addrspace 64)",
      "Context:",
      "96 %56:_(s16), %33:_(s64) :: (store (s16) into %ir.22, align 4, addrspace 64)",
      "%56:_(s16) = 179 %55:_(s32)",
      "%33:_(s64) = 80 %0:_(p0)",
      "%55:_(s32) = 165 %53:_, %54:_",
      "%0:_(p0) = 70 %stack.0",
      "%53:_(s32) = 136 %51:_(s1), %48:_, %52:_",
      "%54:_(s32) = 116 intrinsic(@llvm.agx2.sinc), %53:_(s32)",
      "%51:_(s1) = 134 intpred(eq), %50:_(s32), %9:_",
      "%48:_(s32) = 136 %46:_(s1), %44:_, %47:_",
      "%52:_(s32) = 177 %48:_",
      "%50:_(s32) = 65 %37:_, %49:_",
      "%9:_(s32) = 120 i32 0",
      "%46:_(s1) = 134 intpred(eq), %45:_(s32), %9:_",
      "%44:_(s32) = 164 %43:_, %42:_",
      "%47:_(s32) = 177 %41:_",
      "%37:_(s32) = 166 %34:_, %35:_, %36:_",
      "%49:_(s32) = 120 i32 2",
      "%45:_(s32) = 65 %37:_, %7:_",
      "%43:_(s32) = 121 float 1.000000e+00"
    ]
  },

[maleadt]

Reduced:

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-n8:16:32"
target triple = "air64-apple-macosx15.3.0"

declare half @air.sincos.f16(half, i64)

define void @_Z10intr_test314MtlDeviceArrayI7Float16Li1ELi1EES1_({ i8 addrspace(1)*, [1 x i64] } addrspace(1)* %0, { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %1, i32 %thread_position_in_grid) {
conversion:
  %2 = sext i32 %thread_position_in_grid to i64
  %3 = getelementptr half, half addrspace(1)* null, i64 %2
  %4 = load half, half addrspace(1)* %3, align 2
  %5 = call half @air.sincos.f16(half %4, i64 0)
  ret void
}

!air.kernel = !{!0}
!air.version = !{!6}

!0 = !{void ({ i8 addrspace(1)*, [1 x i64] } addrspace(1)*, { i8 addrspace(1)*, [1 x i64] } addrspace(1)*, i32)* @_Z10intr_test314MtlDeviceArrayI7Float16Li1ELi1EES1_, !1, !2}
!1 = !{}
!2 = !{!3, !4, !5}
!3 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 16, !"air.arg_type_align_size", i32 8, !"air.arg_type_name", !"MtlDeviceVector{Float16, 1}", !"air.arg_name", !"arr_sin"}
!4 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 16, !"air.arg_type_align_size", i32 8, !"air.arg_type_name", !"MtlDeviceVector{Float16, 1}", !"air.arg_name", !"arr_cos"}
!5 = !{i32 2, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint"}
!6 = !{i32 2, i32 5, i32 0}

---

FWIW, the way I reduce these: either run ./metallib-as INPUT-o - | ./metallib-load in a loop, watching for Compiler encountered an internal error, and checking the IPS files in ~/Library/Logs/DiagnosticReports for the expected error condition (here, a legalization error involving @llvm.agx2.sinc); or run the metal driver directly:

❯ cat broken.mtlp-json
{
  "pipelines": {
    "compute_pipelines": [
      {
        "compute_function": "_Z10intr_test314MtlDeviceArrayI7Float16Li1ELi1EES1_",
        "threadgroup_size_is_multiple_of_thread_execution_width": true,
        "max_total_threads_per_threadgroup": 704,
      },
    ]
  }
}

❯ xcrun metal-arch
applegpu_g15s

❯ ./metallib-as broken.ll -f

❯ xcrun metal-tt -arch applegpu_g15s -o main.gpubin broken.mtlp-json broken.metallib
LLVM ERROR: unable to legalize instruction: 96 %56:_(s16), %33:_(s64) :: (store (s16) into %ir.22, align 4, addrspace 64)
Context:
96 %56:_(s16), %33:_(s64) :: (store (s16) into %ir.22, align 4, addrspace 64)
%56:_(s16) = 179 %55:_(s32)
%33:_(s64) = 80 %0:_(p0)
%55:_(s32) = 165 %53:_, %54:_
%0:_(p0) = 70 %stack.0
%53:_(s32) = 136 %51:_(s1), %48:_, %52:_
%54:_(s32) = 116 intrinsic(@llvm.agx2.sinc), %53:_(s32)
%51:_(s1) = 134 intpred(eq), %50:_(s32), %9:_
%48:_(s32) = 136 %46:_(s1), %44:_, %47:_
%52:_(s32) = 177 %48:_
%50:_(s32) = 65 %37:_, %49:_
%9:_(s32) = 120 i32 0
%46:_(s1) = 134 intpred(eq), %45:_(s32), %9:_
%44:_(s32) = 164 %43:_, %42:_
%47:_(s32) = 177 %41:_
%37:_(s32) = 166 %34:_, %35:_, %36:_
%49:_(s32) = 120 i32 2
%45:_(s32) = 65 %37:_, %7:_
%43:_(s32) = 121 float 1.000000e+00
%42:_(s32) = 184 %41:_
 (in function: agc.main)
metal-tt: applegpu-nt command failed

Then I run the above under llvm-reduce.

[maleadt]

declare half @air.sincos.f16(half, half* nocapture writeonly)

vs

declare half @air.sincos.f16(half, i64)

So the problem is we aren't passing a pointer, but an i64. This can be fixed by using @typed_ccall from LLVM.jl instead of plain ccall.

It's also probably useful to add those annotations to https://github.com/JuliaGPU/GPUCompiler.jl/blob/47a4a8b5f6f6737628b2307e76203da262f164a6/src/metal.jl#L1161-L1203

[christiangnrd]

That worked! I'll submit a PR soon. Interesting that it worked with float but not half.

Also, for myself and others reading this in the future, I had to remove the extern in front of the air intrinsic for it to work. What does extern do exactly? I read that in the context of a function call it means that it's been declared elsewhere? Why would I have to remove it when using @typed_ccall? I apologize if this isn't an appropriate place to ask.

[maleadt]

The extern is a way to convince Julia's ccall to accept whatever you throw at it; normally it only allows calling known LLVM intrinsincs (when something starts with a @). But @typed_ccall does not use ccall, it generates direct llvmcall calls, so that isn't required.