; LLVM IR from moarseniy/CudaFEMproject :: FEMproject/matrix_pack/src/cu/cuda_matrix.ll (ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf)

; Unoptimized IR (clang -disable-llvm-passes):
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

%class.Matrix = type <{ ptr, i32, [4 x i8], i64, i64, i64, ptr, i8, [7 x i8] }>

@_ZN11CUDA_Matrix7_handleE = external dso_local global ptr, align 8
@.str = external hidden unnamed_addr constant [104 x i8], align 1
@__FUNCTION__._ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf = external hidden unnamed_addr constant [4 x i8], align 1

declare void @_Z10_my_assertbPKciS0_(i1 noundef zeroext, ptr noundef, i32 noundef, ptr noundef) #0

declare noundef ptr @_ZNK6Matrix8get_dataEv(ptr noundef nonnull align 8 dereferenceable(49)) #0

declare noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49)) #0

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

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

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf(ptr noundef nonnull align 8 dereferenceable(49) %0, ptr noundef nonnull align 8 dereferenceable(49) %1, i64 noundef %2, i64 noundef %3, i1 noundef zeroext %4, ptr noundef nonnull align 8 dereferenceable(49) %5, i64 noundef %6, i1 noundef zeroext %7, i64 noundef %8, float noundef %9) unnamed_addr #2 align 2 {
  %11 = alloca ptr, align 8
  %12 = alloca ptr, align 8
  %13 = alloca i64, align 8
  %14 = alloca i64, align 8
  %15 = alloca i8, align 1
  %16 = alloca ptr, align 8
  %17 = alloca i64, align 8
  %18 = alloca i8, align 1
  %19 = alloca i64, align 8
  %20 = alloca float, align 4
  %21 = alloca float, align 4
  %22 = alloca i32, align 4
  %23 = alloca i32, align 4
  %24 = alloca i32, align 4
  store ptr %0, ptr %11, align 8, !tbaa !5
  store ptr %1, ptr %12, align 8, !tbaa !5
  store i64 %2, ptr %13, align 8, !tbaa !9
  store i64 %3, ptr %14, align 8, !tbaa !9
  %25 = zext i1 %4 to i8
  store i8 %25, ptr %15, align 1, !tbaa !11
  store ptr %5, ptr %16, align 8, !tbaa !5
  store i64 %6, ptr %17, align 8, !tbaa !9
  %26 = zext i1 %7 to i8
  store i8 %26, ptr %18, align 1, !tbaa !11
  store i64 %8, ptr %19, align 8, !tbaa !9
  store float %9, ptr %20, align 4, !tbaa !13
  %27 = load ptr, ptr %11, align 8
  %28 = load ptr, ptr %12, align 8, !tbaa !5
  %29 = call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %28)
  %30 = icmp ugt i64 %29, 0
  br i1 %30, label %31, label %38

31:                                               ; preds = %10
  %32 = load ptr, ptr %16, align 8, !tbaa !5
  %33 = call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %32)
  %34 = icmp ugt i64 %33, 0
  br i1 %34, label %35, label %38

35:                                               ; preds = %31
  %36 = load i64, ptr %19, align 8, !tbaa !9
  %37 = icmp ugt i64 %36, 0
  br label %38

38:                                               ; preds = %35, %31, %10
  %39 = phi i1 [ false, %31 ], [ false, %10 ], [ %37, %35 ]
  call void @_Z10_my_assertbPKciS0_(i1 noundef zeroext %39, ptr noundef @.str, i32 noundef 254, ptr noundef @__FUNCTION__._ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf)
  call void @llvm.lifetime.start.p0(i64 4, ptr %21) #3
  store float 0.000000e+00, ptr %21, align 4, !tbaa !13
  call void @llvm.lifetime.start.p0(i64 4, ptr %22) #3
  %40 = load i64, ptr %13, align 8, !tbaa !9
  %41 = load i64, ptr %14, align 8, !tbaa !9
  %42 = mul i64 %40, %41
  %43 = load ptr, ptr %12, align 8, !tbaa !5
  %44 = call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %43)
  %45 = icmp eq i64 %42, %44
  br i1 %45, label %46, label %47

46:                                               ; preds = %38
  br label %51

47:                                               ; preds = %38
  %48 = load i64, ptr %13, align 8, !tbaa !9
  %49 = load i64, ptr %14, align 8, !tbaa !9
  %50 = mul i64 %48, %49
  br label %51

51:                                               ; preds = %47, %46
  %52 = phi i64 [ 0, %46 ], [ %50, %47 ]
  %53 = trunc i64 %52 to i32
  store i32 %53, ptr %22, align 4, !tbaa !15
  call void @llvm.lifetime.start.p0(i64 4, ptr %23) #3
  %54 = load i64, ptr %14, align 8, !tbaa !9
  %55 = load i64, ptr %17, align 8, !tbaa !9
  %56 = mul i64 %54, %55
  %57 = load ptr, ptr %16, align 8, !tbaa !5
  %58 = call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %57)
  %59 = icmp eq i64 %56, %58
  br i1 %59, label %60, label %61

60:                                               ; preds = %51
  br label %65

61:                                               ; preds = %51
  %62 = load i64, ptr %14, align 8, !tbaa !9
  %63 = load i64, ptr %17, align 8, !tbaa !9
  %64 = mul i64 %62, %63
  br label %65

65:                                               ; preds = %61, %60
  %66 = phi i64 [ 0, %60 ], [ %64, %61 ]
  %67 = trunc i64 %66 to i32
  store i32 %67, ptr %23, align 4, !tbaa !15
  call void @llvm.lifetime.start.p0(i64 4, ptr %24) #3
  %68 = load i64, ptr %13, align 8, !tbaa !9
  %69 = load i64, ptr %17, align 8, !tbaa !9
  %70 = mul i64 %68, %69
  %71 = trunc i64 %70 to i32
  store i32 %71, ptr %24, align 4, !tbaa !15
  %72 = load ptr, ptr @_ZN11CUDA_Matrix7_handleE, align 8, !tbaa !5
  %73 = load i8, ptr %15, align 1, !tbaa !11, !range !17, !noundef !18
  %74 = trunc i8 %73 to i1
  %75 = zext i1 %74 to i64
  %76 = select i1 %74, i32 1, i32 0
  %77 = load i8, ptr %18, align 1, !tbaa !11, !range !17, !noundef !18
  %78 = trunc i8 %77 to i1
  %79 = zext i1 %78 to i64
  %80 = select i1 %78, i32 1, i32 0
  %81 = load i64, ptr %13, align 8, !tbaa !9
  %82 = trunc i64 %81 to i32
  %83 = load i64, ptr %17, align 8, !tbaa !9
  %84 = trunc i64 %83 to i32
  %85 = load i64, ptr %14, align 8, !tbaa !9
  %86 = trunc i64 %85 to i32
  %87 = load ptr, ptr %12, align 8, !tbaa !5
  %88 = call noundef ptr @_ZNK6Matrix8get_dataEv(ptr noundef nonnull align 8 dereferenceable(49) %87)
  %89 = load i8, ptr %15, align 1, !tbaa !11, !range !17, !noundef !18
  %90 = trunc i8 %89 to i1
  br i1 %90, label %91, label %93

91:                                               ; preds = %65
  %92 = load i64, ptr %14, align 8, !tbaa !9
  br label %95

93:                                               ; preds = %65
  %94 = load i64, ptr %13, align 8, !tbaa !9
  br label %95

95:                                               ; preds = %93, %91
  %96 = phi i64 [ %92, %91 ], [ %94, %93 ]
  %97 = trunc i64 %96 to i32
  %98 = load i32, ptr %22, align 4, !tbaa !15
  %99 = sext i32 %98 to i64
  %100 = load ptr, ptr %16, align 8, !tbaa !5
  %101 = call noundef ptr @_ZNK6Matrix8get_dataEv(ptr noundef nonnull align 8 dereferenceable(49) %100)
  %102 = load i8, ptr %18, align 1, !tbaa !11, !range !17, !noundef !18
  %103 = trunc i8 %102 to i1
  br i1 %103, label %104, label %106

104:                                              ; preds = %95
  %105 = load i64, ptr %17, align 8, !tbaa !9
  br label %108

106:                                              ; preds = %95
  %107 = load i64, ptr %14, align 8, !tbaa !9
  br label %108

108:                                              ; preds = %106, %104
  %109 = phi i64 [ %105, %104 ], [ %107, %106 ]
  %110 = trunc i64 %109 to i32
  %111 = load i32, ptr %23, align 4, !tbaa !15
  %112 = sext i32 %111 to i64
  %113 = getelementptr inbounds %class.Matrix, ptr %27, i32 0, i32 6
  %114 = load ptr, ptr %113, align 8, !tbaa !19
  %115 = load i64, ptr %13, align 8, !tbaa !9
  %116 = trunc i64 %115 to i32
  %117 = load i32, ptr %24, align 4, !tbaa !15
  %118 = sext i32 %117 to i64
  %119 = load i64, ptr %19, align 8, !tbaa !9
  %120 = trunc i64 %119 to i32
  %121 = call i32 @cublasSgemmStridedBatched(ptr noundef %72, i32 noundef %76, i32 noundef %80, i32 noundef %82, i32 noundef %84, i32 noundef %86, ptr noundef %20, ptr noundef %88, i32 noundef %97, i64 noundef %99, ptr noundef %101, i32 noundef %110, i64 noundef %112, ptr noundef %21, ptr noundef %114, i32 noundef %116, i64 noundef %118, i32 noundef %120)
  %122 = call i32 @cudaDeviceSynchronize()
  call void @llvm.lifetime.end.p0(i64 4, ptr %24) #3
  call void @llvm.lifetime.end.p0(i64 4, ptr %23) #3
  call void @llvm.lifetime.end.p0(i64 4, ptr %22) #3
  call void @llvm.lifetime.end.p0(i64 4, ptr %21) #3
  ret void
}

declare i32 @cublasSgemmStridedBatched(ptr noundef, i32 noundef, i32 noundef, i32 noundef, i32 noundef, i32 noundef, ptr noundef, ptr noundef, i32 noundef, i64 noundef, ptr noundef, i32 noundef, i64 noundef, ptr noundef, ptr noundef, i32 noundef, i64 noundef, i32 noundef) #0

declare i32 @cudaDeviceSynchronize() #0

attributes #0 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #1 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #2 = { mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { nounwind }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 8, !"PIC Level", i32 2}
!2 = !{i32 7, !"PIE Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{!"clang version 19.1.0 (git@github.com:llvm/llvm-project.git a4bf6cd7cfb1a1421ba92bca9d017b49936c55e4)"}
!5 = !{!6, !6, i64 0}
!6 = !{!"any pointer", !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!10, !10, i64 0}
!10 = !{!"long", !7, i64 0}
!11 = !{!12, !12, i64 0}
!12 = !{!"bool", !7, i64 0}
!13 = !{!14, !14, i64 0}
!14 = !{!"float", !7, i64 0}
!15 = !{!16, !16, i64 0}
!16 = !{!"int", !7, i64 0}
!17 = !{i8 0, i8 2}
!18 = !{}
!19 = !{!20, !6, i64 40}
!20 = !{!"_ZTS6Matrix", !21, i64 8, !10, i64 16, !10, i64 24, !10, i64 32, !6, i64 40, !12, i64 48}
!21 = !{!"_ZTS11DEVICE_NAME", !7, i64 0}


; Optimized IR (opt -passes='default<O3>'):
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

@_ZN11CUDA_Matrix7_handleE = external dso_local local_unnamed_addr global ptr, align 8
@.str = external hidden unnamed_addr constant [104 x i8], align 1
@__FUNCTION__._ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf = external hidden unnamed_addr constant [4 x i8], align 1

declare void @_Z10_my_assertbPKciS0_(i1 noundef zeroext, ptr noundef, i32 noundef, ptr noundef) local_unnamed_addr #0

declare noundef ptr @_ZNK6Matrix8get_dataEv(ptr noundef nonnull align 8 dereferenceable(49)) local_unnamed_addr #0

declare noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49)) local_unnamed_addr #0

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

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

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf(ptr nocapture noundef nonnull readonly align 8 dereferenceable(49) %0, ptr noundef nonnull align 8 dereferenceable(49) %1, i64 noundef %2, i64 noundef %3, i1 noundef zeroext %4, ptr noundef nonnull align 8 dereferenceable(49) %5, i64 noundef %6, i1 noundef zeroext %7, i64 noundef %8, float noundef %9) unnamed_addr #2 align 2 {
  %11 = alloca float, align 4
  %12 = alloca float, align 4
  store float %9, ptr %11, align 4, !tbaa !5
  %13 = tail call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %1)
  %.not = icmp eq i64 %13, 0
  br i1 %.not, label %17, label %14

14:                                               ; preds = %10
  %15 = tail call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %5)
  %.not33 = icmp ne i64 %15, 0
  %16 = icmp ne i64 %8, 0
  %spec.select = and i1 %16, %.not33
  br label %17

17:                                               ; preds = %14, %10
  %18 = phi i1 [ false, %10 ], [ %spec.select, %14 ]
  tail call void @_Z10_my_assertbPKciS0_(i1 noundef zeroext %18, ptr noundef nonnull @.str, i32 noundef 254, ptr noundef nonnull @__FUNCTION__._ZN11CUDA_Matrix3bmmERK6MatrixmmbS2_mbmf)
  call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %12) #3
  store float 0.000000e+00, ptr %12, align 4, !tbaa !5
  %19 = mul i64 %3, %2
  %20 = tail call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %1)
  %21 = icmp eq i64 %19, %20
  %sext = shl i64 %19, 32
  %22 = ashr exact i64 %sext, 32
  %23 = mul i64 %6, %3
  %24 = tail call noundef i64 @_ZNK6Matrix15get_numElementsEv(ptr noundef nonnull align 8 dereferenceable(49) %5)
  %25 = icmp eq i64 %23, %24
  %sext34 = shl i64 %23, 32
  %26 = ashr exact i64 %sext34, 32
  %27 = load ptr, ptr @_ZN11CUDA_Matrix7_handleE, align 8, !tbaa !9
  %28 = zext i1 %4 to i32
  %29 = zext i1 %7 to i32
  %30 = trunc i64 %2 to i32
  %31 = trunc i64 %6 to i32
  %32 = trunc i64 %3 to i32
  %33 = tail call noundef ptr @_ZNK6Matrix8get_dataEv(ptr noundef nonnull align 8 dereferenceable(49) %1)
  %34 = select i1 %4, i64 %3, i64 %2
  %35 = trunc i64 %34 to i32
  %36 = select i1 %21, i64 0, i64 %22
  %37 = tail call noundef ptr @_ZNK6Matrix8get_dataEv(ptr noundef nonnull align 8 dereferenceable(49) %5)
  %38 = select i1 %7, i64 %6, i64 %3
  %39 = trunc i64 %38 to i32
  %40 = select i1 %25, i64 0, i64 %26
  %41 = getelementptr inbounds i8, ptr %0, i64 40
  %42 = load ptr, ptr %41, align 8, !tbaa !11
  %43 = shl i64 %2, 32
  %sext35 = mul i64 %43, %6
  %44 = ashr exact i64 %sext35, 32
  %45 = trunc i64 %8 to i32
  %46 = call i32 @cublasSgemmStridedBatched(ptr noundef %27, i32 noundef %28, i32 noundef %29, i32 noundef %30, i32 noundef %31, i32 noundef %32, ptr noundef nonnull %11, ptr noundef %33, i32 noundef %35, i64 noundef %36, ptr noundef %37, i32 noundef %39, i64 noundef %40, ptr noundef nonnull %12, ptr noundef %42, i32 noundef %30, i64 noundef %44, i32 noundef %45)
  %47 = call i32 @cudaDeviceSynchronize()
  call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %12) #3
  ret void
}

declare i32 @cublasSgemmStridedBatched(ptr noundef, i32 noundef, i32 noundef, i32 noundef, i32 noundef, i32 noundef, ptr noundef, ptr noundef, i32 noundef, i64 noundef, ptr noundef, i32 noundef, i64 noundef, ptr noundef, ptr noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #0

declare i32 @cudaDeviceSynchronize() local_unnamed_addr #0

attributes #0 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #1 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #2 = { mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { nounwind }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 8, !"PIC Level", i32 2}
!2 = !{i32 7, !"PIE Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{!"clang version 19.1.0 (git@github.com:llvm/llvm-project.git a4bf6cd7cfb1a1421ba92bca9d017b49936c55e4)"}
!5 = !{!6, !6, i64 0}
!6 = !{!"float", !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!10, !10, i64 0}
!10 = !{!"any pointer", !7, i64 0}
!11 = !{!12, !10, i64 40}
!12 = !{!"_ZTS6Matrix", !13, i64 8, !14, i64 16, !14, i64 24, !14, i64 32, !10, i64 40, !15, i64 48}
!13 = !{!"_ZTS11DEVICE_NAME", !7, i64 0}
!14 = !{!"long", !7, i64 0}
!15 = !{!"bool", !7, i64 0}


; Effective O3 passes: EarlyCSEPass,GlobalOptPass,InferFunctionAttrsPass,InstCombinePass,PostOrderFunctionAttrsPass,ReassociatePass,SROAPass,SimplifyCFGPass,TailCallElimPass