bug.ll

; ModuleID = 'module__reduce_dim_1_const_dispatch_0_embedded_elf_x86_64.codegen.bc'
source_filename = "_reduce_dim_1_const_dispatch_0"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-unknown-eabi-elf"

%iree_hal_executable_library_header_t = type { i32, ptr, i32, i32 }
%iree_hal_executable_dispatch_attrs_v0_t = type { i16, i16 }
%iree_hal_executable_src_loc_v0_t = type { i32, i32, ptr }
%iree_hal_executable_library_v0_t = type { ptr, %iree_hal_executable_import_table_v0_t, %iree_hal_executable_export_table_v0_t, %iree_hal_executable_constant_table_v0_t }
%iree_hal_executable_import_table_v0_t = type { i32, ptr }
%iree_hal_executable_export_table_v0_t = type { i32, ptr, ptr, ptr, ptr, ptr }
%iree_hal_executable_constant_table_v0_t = type { i32 }
%iree_hal_executable_dispatch_state_v0_t = type { i32, i32, i16, i16, i32, i32, i16, i8, i8, ptr, ptr, ptr }

@0 = private constant [31 x i8] c"_reduce_dim_1_const_dispatch_0\00", align 1
@iree_hal_executable_library_query_v0_header = private constant %iree_hal_executable_library_header_t { i32 3, ptr @0, i32 0, i32 0 }
@iree_hal_executable_library_query_v0_funcs = private constant [1 x ptr] [ptr @_reduce_dim_1_const_dispatch_0_generic_2x5_i32]
@iree_hal_executable_library_query_v0_attrs = private constant [1 x %iree_hal_executable_dispatch_attrs_v0_t] zeroinitializer
@1 = private constant [47 x i8] c"_reduce_dim_1_const_dispatch_0_generic_2x5_i32\00", align 1
@iree_hal_executable_library_query_v0_names = private constant [1 x ptr] [ptr @1]
@2 = private constant [1 x i8] zeroinitializer, align 1
@iree_hal_executable_library_query_v0_tags = private constant [1 x ptr] [ptr @2]
@3 = private constant [13 x i8] c"reduced.mlir\00", align 1
@iree_hal_executable_library_query_v0_src_locs = private constant [1 x %iree_hal_executable_src_loc_v0_t] [%iree_hal_executable_src_loc_v0_t { i32 5, i32 12, ptr @3 }]
@iree_hal_executable_library_query_v0 = private constant %iree_hal_executable_library_v0_t { ptr @iree_hal_executable_library_query_v0_header, %iree_hal_executable_import_table_v0_t zeroinitializer, %iree_hal_executable_export_table_v0_t { i32 1, ptr @iree_hal_executable_library_query_v0_funcs, ptr @iree_hal_executable_library_query_v0_attrs, ptr @iree_hal_executable_library_query_v0_names, ptr @iree_hal_executable_library_query_v0_tags, ptr @iree_hal_executable_library_query_v0_src_locs }, %iree_hal_executable_constant_table_v0_t zeroinitializer }

declare ptr @malloc(i64) #0

declare void @free(ptr) #0

define internal i32 @_reduce_dim_1_const_dispatch_0_generic_2x5_i32(ptr noalias nonnull align 16 %0, ptr noalias nonnull align 16 %1, ptr noalias nonnull align 16 %2) #0 !dbg !3 {
  %4 = load %iree_hal_executable_dispatch_state_v0_t, ptr %1, align 8, !dbg !79
  %5 = extractvalue %iree_hal_executable_dispatch_state_v0_t %4, 10, !dbg !79
  %6 = load ptr, ptr %5, align 8, !dbg !79
  %7 = ptrtoint ptr %6 to i64, !dbg !79
  %8 = and i64 %7, 63, !dbg !79
  %9 = icmp eq i64 %8, 0, !dbg !79
  call void @llvm.assume(i1 %9), !dbg !79
  %10 = load %iree_hal_executable_dispatch_state_v0_t, ptr %1, align 8, !dbg !80
  %11 = extractvalue %iree_hal_executable_dispatch_state_v0_t %10, 10, !dbg !80
  %12 = getelementptr ptr, ptr %11, i32 1, !dbg !80
  %13 = load ptr, ptr %12, align 8, !dbg !80
  %14 = ptrtoint ptr %13 to i64, !dbg !80
  %15 = and i64 %14, 63, !dbg !80
  %16 = icmp eq i64 %15, 0, !dbg !80
  call void @llvm.assume(i1 %16), !dbg !80
  store <2 x i32> <i32 10, i32 10>, ptr %13, align 4, !dbg !81
  br label %17, !dbg !82

17:                                               ; preds = %47, %3
  %18 = phi i64 [ %48, %47 ], [ 0, %3 ]
  %19 = icmp slt i64 %18, 2, !dbg !82
  br i1 %19, label %20, label %49, !dbg !82

20:                                               ; preds = %17
  %21 = getelementptr i32, ptr %13, i64 %18, !dbg !82
  %22 = call <2 x i32> @llvm.masked.load.v2i32.p0(ptr %21, i32 4, <2 x i1> <i1 true, i1 false>, <2 x i32> zeroinitializer), !dbg !82
  br label %23, !dbg !82

23:                                               ; preds = %27, %20
  %24 = phi i64 [ %46, %27 ], [ 0, %20 ]
  %25 = phi <2 x i32> [ %45, %27 ], [ %22, %20 ]
  %26 = icmp slt i64 %24, 5, !dbg !82
  br i1 %26, label %27, label %47, !dbg !82

27:                                               ; preds = %23
  %28 = mul i64 %24, -1, !dbg !82
  %29 = add i64 %28, 5, !dbg !82
  %30 = icmp slt i64 %29, 4, !dbg !82
  %31 = select i1 %30, i64 %29, i64 4, !dbg !82
  %32 = mul i64 %18, 5, !dbg !82
  %33 = add i64 %32, %24, !dbg !82
  %34 = insertelement <4 x i64> undef, i64 %31, i32 0, !dbg !82
  %35 = shufflevector <4 x i64> %34, <4 x i64> undef, <4 x i32> zeroinitializer, !dbg !82
  %36 = icmp sgt <4 x i64> %35, <i64 0, i64 1, i64 2, i64 3>, !dbg !82
  %37 = getelementptr i32, ptr %6, i64 %33, !dbg !82
  %38 = getelementptr i32, ptr %37, i64 0, !dbg !82
  %39 = call <4 x i32> @llvm.masked.load.v4i32.p0(ptr %38, i32 4, <4 x i1> %36, <4 x i32> zeroinitializer), !dbg !82
  %40 = extractelement <2 x i32> %25, i64 0, !dbg !83
  %41 = call i32 @llvm.vp.reduce.add.v4i32(i32 %40, <4 x i32> %39, <4 x i1> %36, i32 4), !dbg !83
  %42 = insertelement <2 x i32> zeroinitializer, i32 %41, i64 0, !dbg !83
  %43 = extractelement <2 x i32> %25, i64 1, !dbg !83
  %44 = call i32 @llvm.vp.reduce.add.v4i32(i32 %43, <4 x i32> zeroinitializer, <4 x i1> zeroinitializer, i32 4), !dbg !83
  %45 = insertelement <2 x i32> %42, i32 %44, i64 1, !dbg !83
  %46 = add i64 %24, 4, !dbg !82
  br label %23, !dbg !82

47:                                               ; preds = %23
  call void @llvm.masked.store.v2i32.p0(<2 x i32> %25, ptr %21, i32 4, <2 x i1> <i1 true, i1 false>), !dbg !83
  %48 = add i64 %18, 1, !dbg !82
  br label %17, !dbg !82

49:                                               ; preds = %17
  ret i32 0, !dbg !84
}

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite)
declare void @llvm.assume(i1 noundef) #1

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: read)
declare <2 x i32> @llvm.masked.load.v2i32.p0(ptr nocapture, i32 immarg, <2 x i1>, <2 x i32>) #2

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: write)
declare void @llvm.masked.store.v2i32.p0(<2 x i32>, ptr nocapture, i32 immarg, <2 x i1>) #3

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: read)
declare <4 x i32> @llvm.masked.load.v4i32.p0(ptr nocapture, i32 immarg, <4 x i1>, <4 x i32>) #2

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare i32 @llvm.vp.reduce.add.v4i32(i32, <4 x i32>, <4 x i1>, i32) #4

; Function Attrs: uwtable
define dso_local dllexport ptr @iree_hal_executable_library_query(i32 %0, ptr %1) #5 {
entry:
  %2 = icmp eq i32 %0, 3
  %3 = select i1 %2, ptr @iree_hal_executable_library_query_v0, ptr null
  ret ptr %3
}

attributes #0 = { "frame-pointer"="all" "hot" "no-builtins" "nonlazybind" }
attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) }
attributes #2 = { nocallback nofree nosync nounwind willreturn memory(argmem: read) }
attributes #3 = { nocallback nofree nosync nounwind willreturn memory(argmem: write) }
attributes #4 = { nocallback nofree nosync nounwind willreturn memory(none) }
attributes #5 = { uwtable "nonlazybind" }

!llvm.module.flags = !{!0}
!llvm.dbg.cu = !{!1}

!0 = !{i32 2, !"Debug Info Version", i32 3}
!1 = distinct !DICompileUnit(language: DW_LANG_C17, file: !2, producer: "IREE", isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug)
!2 = !DIFile(filename: "reduced.mlir", directory: "")
!3 = distinct !DISubprogram(name: "_reduce_dim_1_const_dispatch_0_generic_2x5_i32", linkageName: "_reduce_dim_1_const_dispatch_0_generic_2x5_i32", scope: !2, file: !2, line: 1, type: !4, scopeLine: 1, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !1)
!4 = !DISubroutineType(cc: DW_CC_normal, types: !5)
!5 = !{!6, !7, !38, !67}
!6 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
!7 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !8, size: 64)
!8 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !9)
!9 = !DIDerivedType(tag: DW_TAG_typedef, name: "iree_hal_executable_environment_v0_t", baseType: !10)
!10 = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "iree_hal_executable_environment_v0_t", scope: !11, file: !11, line: 246, size: 768, elements: !12)
!11 = !DIFile(filename: "runtime/src/iree/hal/local/executable_library.h", directory: ".")
!12 = !{!13, !21, !24, !27, !29}
!13 = !DIDerivedType(tag: DW_TAG_member, name: "constants", baseType: !14, size: 64)
!14 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !15, size: 64)
!15 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !16)
!16 = !DICompositeType(tag: DW_TAG_array_type, scope: !11, file: !11, line: 227, baseType: !17, size: 2048, elements: !19)
!17 = !DIDerivedType(tag: DW_TAG_typedef, name: "uint32_t", baseType: !18)
!18 = !DIBasicType(name: "unsigned int", size: 32, encoding: DW_ATE_unsigned)
!19 = !{!20}
!20 = !DISubrange(count: 64)
!21 = !DIDerivedType(tag: DW_TAG_member, name: "import_thunk", baseType: !22, size: 64, offset: 64)
!22 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !23, size: 64)
!23 = !DIBasicType(name: "void", encoding: DW_ATE_address)
!24 = !DIDerivedType(tag: DW_TAG_member, name: "import_funcs", baseType: !25, size: 64, offset: 128)
!25 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !26, size: 64)
!26 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !22)
!27 = !DIDerivedType(tag: DW_TAG_member, name: "import_contexts", baseType: !28, size: 64, offset: 192)
!28 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !25, size: 64)
!29 = !DIDerivedType(tag: DW_TAG_member, name: "processor", baseType: !30, offset: 256)
!30 = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "iree_hal_processor_v0_t", scope: !11, file: !11, line: 227, size: 512, elements: !31)
!31 = !{!32}
!32 = !DIDerivedType(tag: DW_TAG_member, name: "data", baseType: !33)
!33 = !DICompositeType(tag: DW_TAG_array_type, scope: !11, file: !11, line: 227, baseType: !34, size: 512, elements: !36)
!34 = !DIDerivedType(tag: DW_TAG_typedef, name: "uint64_t", baseType: !35)
!35 = !DIBasicType(name: "long long unsigned int", size: 64, encoding: DW_ATE_unsigned)
!36 = !{!37}
!37 = !DISubrange(count: 8)
!38 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !39, size: 64)
!39 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !40)
!40 = !DIDerivedType(tag: DW_TAG_typedef, name: "iree_hal_executable_dispatch_state_v0_t", baseType: !41)
!41 = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "iree_hal_executable_dispatch_state_v0_t", scope: !11, file: !11, line: 275, size: 384, elements: !42)
!42 = !{!43, !44, !45, !48, !49, !50, !51, !52, !55, !56, !57, !62}
!43 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_size_x", baseType: !17, size: 32)
!44 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_size_y", baseType: !17, size: 32, offset: 32)
!45 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_size_z", baseType: !46, size: 16, offset: 64)
!46 = !DIDerivedType(tag: DW_TAG_typedef, name: "uint16_t", baseType: !47)
!47 = !DIBasicType(name: "unsigned short", size: 16, encoding: DW_ATE_unsigned)
!48 = !DIDerivedType(tag: DW_TAG_member, name: "push_constant_count", baseType: !46, size: 16, offset: 80)
!49 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_count_x", baseType: !17, size: 32, offset: 96)
!50 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_count_y", baseType: !17, size: 32, offset: 128)
!51 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_count_z", baseType: !46, size: 16, offset: 160)
!52 = !DIDerivedType(tag: DW_TAG_member, name: "max_concurrency", baseType: !53, size: 8, offset: 176)
!53 = !DIDerivedType(tag: DW_TAG_typedef, name: "uint8_t", baseType: !54)
!54 = !DIBasicType(name: "unsigned char", size: 8, encoding: DW_ATE_unsigned_char)
!55 = !DIDerivedType(tag: DW_TAG_member, name: "binding_count", baseType: !53, size: 8, offset: 184)
!56 = !DIDerivedType(tag: DW_TAG_member, name: "push_constants", baseType: !14, size: 64, offset: 192)
!57 = !DIDerivedType(tag: DW_TAG_member, name: "binding_ptrs", baseType: !58, size: 64, offset: 256)
!58 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !59, size: 64)
!59 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !60)
!60 = !DICompositeType(tag: DW_TAG_array_type, scope: !11, file: !11, line: 227, baseType: !61, size: 4096, elements: !19)
!61 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !53, size: 64)
!62 = !DIDerivedType(tag: DW_TAG_member, name: "binding_lengths", baseType: !63, size: 64, offset: 320)
!63 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !64, size: 64)
!64 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !65)
!65 = !DICompositeType(tag: DW_TAG_array_type, scope: !11, file: !11, line: 227, baseType: !66, size: 4096, elements: !19)
!66 = !DIDerivedType(tag: DW_TAG_typedef, name: "size_t", baseType: !34)
!67 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !68, size: 64)
!68 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !69)
!69 = !DIDerivedType(tag: DW_TAG_typedef, name: "iree_hal_executable_workgroup_state_v0_t", baseType: !70)
!70 = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "iree_hal_executable_workgroup_state_v0_t", scope: !11, file: !11, line: 321, size: 256, elements: !71)
!71 = !{!72, !73, !74, !75, !76, !77, !78}
!72 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_id_x", baseType: !17, size: 32)
!73 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_id_y", baseType: !17, size: 32, offset: 32)
!74 = !DIDerivedType(tag: DW_TAG_member, name: "workgroup_id_z", baseType: !46, size: 16, offset: 64)
!75 = !DIDerivedType(tag: DW_TAG_member, name: "reserved", baseType: !46, size: 16, offset: 80)
!76 = !DIDerivedType(tag: DW_TAG_member, name: "processor_id", baseType: !17, size: 32, offset: 96)
!77 = !DIDerivedType(tag: DW_TAG_member, name: "local_memory", baseType: !22, size: 64, offset: 128)
!78 = !DIDerivedType(tag: DW_TAG_member, name: "local_memory_size", baseType: !17, size: 32, offset: 192)
!79 = !DILocation(line: 10, column: 65, scope: !3)
!80 = !DILocation(line: 10, column: 121, scope: !3)
!81 = !DILocation(line: 11, column: 20, scope: !3)
!82 = !DILocation(line: 15, column: 14, scope: !3)
!83 = !DILocation(line: 17, column: 16, scope: !3)
!84 = !DILocation(line: 21, column: 9, scope: !3)