Skip to content

Instantly share code, notes, and snippets.

@makslevental

makslevental/vector_add.air.ll

Created November 22, 2025 18:57
Show Gist options

  • Select an option

    Learn more about clone URLs
  • Save makslevental/8e115ba84489d32bf5c043662a1ba068 to your computer and use it in GitHub Desktop.

Select an option

Learn more about clone URLs
Save makslevental/8e115ba84489d32bf5c043662a1ba068 to your computer and use it in GitHub Desktop.
Download ZIP
Raw
vector_add.air.ll
; ModuleID = 'vector_add.air'
source_filename = "vector_add.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_v27-apple-macosx15.6.0"
; Function Attrs: mustprogress nounwind
define weak_odr void @mlir_vector_add(ptr addrspace(1) noundef "air-buffer-no-alias" %0, ptr addrspace(1) noundef "air-buffer-no-alias" %1, ptr addrspace(1) noundef "air-buffer-no-alias" %2, <3 x i32> noundef %3) local_unnamed_addr #0 {
%5 = extractelement <3 x i32> %3, i64 0
%6 = zext i32 %5 to i64
%7 = getelementptr inbounds half, ptr addrspace(1) %0, i64 %6
%8 = load half, ptr addrspace(1) %7, align 2, !tbaa !23, !alias.scope !27, !noalias !30
%9 = getelementptr inbounds half, ptr addrspace(1) %1, i64 %6
%10 = load half, ptr addrspace(1) %9, align 2, !tbaa !23, !alias.scope !33, !noalias !34
%11 = fadd half %8, %10
%12 = fpext half %11 to float
%13 = fadd float %12, 0x40091EB860000000
%14 = fptrunc float %13 to half
%15 = getelementptr inbounds half, ptr addrspace(1) %2, i64 %6
store half %14, ptr addrspace(1) %15, align 2, !tbaa !23, !alias.scope !35, !noalias !36
ret void
}
attributes #0 = { mustprogress nounwind "frame-pointer"="all" "min-legal-vector-width"="96" "no-builtins" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
!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 26, i32 1]}
!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 = !{ptr @mlir_vector_add, !10, !11}
!10 = !{}
!11 = !{!12, !13, !14, !15}
!12 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.read", !"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", !"a"}
!13 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, i32 1, !"air.read", !"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", !"b"}
!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", !"result"}
!15 = !{i32 3, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint3", !"air.arg_name", !"thread_position_in_grid"}
!16 = !{!"air.compile.denorms_disable"}
!17 = !{!"air.compile.fast_math_disable"}
!18 = !{!"air.compile.framebuffer_fetch_enable"}
!19 = !{!"Apple metal version 32023.830 (metalfe-32023.830.2)"}
!20 = !{i32 2, i32 7, i32 0}
!21 = !{!"Metal", i32 3, i32 2, i32 0}
!22 = !{!"/Users/mlevental/dev_projects/mlir-apple-gpu/vector_add.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(0)"}
!29 = distinct !{!29, !"air-alias-scopes(mlir_vector_add)"}
!30 = !{!31, !32}
!31 = distinct !{!31, !29, !"air-alias-scope-arg(1)"}
!32 = distinct !{!32, !29, !"air-alias-scope-arg(2)"}
!33 = !{!31}
!34 = !{!28, !32}
!35 = !{!32}
!36 = !{!28, !31}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment