; ModuleID = 'copy_mat'
source_filename = "copy_mat"
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

@kernel_0 = private unnamed_addr constant [2653 x i8] c"//\0A// Generated by LLVM NVPTX Back-End\0A//\0A\0A.version 3.2\0A.target sm_30\0A.address_size 64\0A\0A\09// .globl\09kernel_0\0A\0A.visible .entry kernel_0(\0A\09.param .u64 kernel_0_param_0,\0A\09.param .u64 kernel_0_param_1,\0A\09.param .u64 kernel_0_param_2,\0A\09.param .u64 kernel_0_param_3,\0A\09.param .u64 kernel_0_param_4\0A)\0A.maxntid 32, 16, 1\0A{\0A\09.reg .pred \09%p<6>;\0A\09.reg .b32 \09%r<5>;\0A\09.reg .b64 \09%rd<74>;\0A\0A\09ld.param.u64 \09%rd32, [kernel_0_param_3];\0A\09ld.param.u64 \09%rd34, [kernel_0_param_0];\0A\09ld.param.u64 \09%rd35, [kernel_0_param_1];\0A\09mov.u32 \09%r1, %ctaid.x;\0A\09ld.param.u64 \09%rd36, [kernel_0_param_2];\0A\09mov.u32 \09%r2, %ctaid.y;\0A\09ld.param.u64 \09%rd37, [kernel_0_param_4];\0A\09mov.u32 \09%r3, %tid.x;\0A\09cvt.u64.u32 \09%rd38, %r3;\0A\09mov.u32 \09%r4, %tid.y;\0A\09cvt.u64.u32 \09%rd39, %r4;\0A\09mul.wide.u32 \09%rd40, %r1, 32;\0A\09add.s64 \09%rd41, %rd32, -1;\0A\09sub.s64 \09%rd42, %rd41, %rd40;\0A\09shr.u64 \09%rd1, %rd42, 13;\0A\09add.s64 \09%rd2, %rd40, %rd38;\0A\09mul.wide.u32 \09%rd43, %r2, 32;\0A\09add.s64 \09%rd44, %rd36, -1;\0A\09sub.s64 \09%rd45, %rd44, %rd43;\0A\09shr.u64 \09%rd3, %rd45, 13;\0A\09mul.wide.u32 \09%rd46, %r2, 2;\0A\09neg.s64 \09%rd4, %rd46;\0A\09sub.s64 \09%rd47, %rd44, %rd39;\0A\09shr.s64 \09%rd5, %rd47, 4;\0A\09add.s64 \09%rd48, %rd43, %rd39;\0A\09mul.lo.s64 \09%rd49, %rd32, %rd48;\0A\09add.s64 \09%rd50, %rd49, %rd40;\0A\09add.s64 \09%rd51, %rd50, %rd38;\0A\09shl.b64 \09%rd52, %rd51, 3;\0A\09add.s64 \09%rd66, %rd35, %rd52;\0A\09shl.b64 \09%rd7, %rd32, 16;\0A\09shl.b64 \09%rd8, %rd32, 7;\0A\09mul.lo.s64 \09%rd53, %rd37, %rd48;\0A\09add.s64 \09%rd54, %rd53, %rd40;\0A\09add.s64 \09%rd55, %rd54, %rd38;\0A\09shl.b64 \09%rd56, %rd55, 3;\0A\09add.s64 \09%rd65, %rd34, %rd56;\0A\09shl.b64 \09%rd10, %rd37, 16;\0A\09shl.b64 \09%rd11, %rd37, 7;\0A\09mov.u64 \09%rd67, 0;\0ALBB0_1:\0A\09shl.b64 \09%rd57, %rd67, 13;\0A\09add.s64 \09%rd58, %rd2, %rd57;\0A\09setp.ge.s64 \09%p1, %rd58, %rd32;\0A\09@%p1 bra \09LBB0_5;\0A\09mov.u64 \09%rd70, 0;\0A\09mov.u64 \09%rd68, %rd65;\0A\09mov.u64 \09%rd69, %rd66;\0ALBB0_3:\0A\09shl.b64 \09%rd60, %rd70, 9;\0A\09sub.s64 \09%rd61, %rd4, %rd60;\0A\09add.s64 \09%rd62, %rd61, %rd5;\0A\09min.s64 \09%rd21, %rd62, 1;\0A\09setp.lt.s64 \09%p2, %rd21, 0;\0A\09@%p2 bra \09LBB0_8;\0A\09add.s64 \09%rd22, %rd21, -1;\0A\09mov.u64 \09%rd71, -1;\0A\09mov.u64 \09%rd72, %rd68;\0A\09mov.u64 \09%rd73, %rd69;\0ALBB0_7:\0A\09ld.global.u64 \09%rd64, [%rd72];\0A\09st.global.u64 \09[%rd73], %rd64;\0A\09add.s64 \09%rd73, %rd73, %rd8;\0A\09add.s64 \09%rd72, %rd72, %rd11;\0A\09add.s64 \09%rd71, %rd71, 1;\0A\09setp.gt.s64 \09%p3, %rd71, %rd22;\0A\09@%p3 bra \09LBB0_8;\0A\09bra.uni \09LBB0_7;\0ALBB0_8:\0A\09add.s64 \09%rd23, %rd70, 1;\0A\09add.s64 \09%rd69, %rd69, %rd7;\0A\09add.s64 \09%rd68, %rd68, %rd10;\0A\09setp.eq.s64 \09%p4, %rd70, %rd3;\0A\09mov.u64 \09%rd70, %rd23;\0A\09@%p4 bra \09LBB0_5;\0A\09bra.uni \09LBB0_3;\0ALBB0_5:\0A\09add.s64 \09%rd15, %rd67, 1;\0A\09add.s64 \09%rd66, %rd66, 65536;\0A\09add.s64 \09%rd65, %rd65, 65536;\0A\09setp.eq.s64 \09%p5, %rd67, %rd1;\0A\09mov.u64 \09%rd67, %rd15;\0A\09@%p5 bra \09LBB0_6;\0A\09bra.uni \09LBB0_1;\0ALBB0_6:\0A\09ret;\0A}\0A\0A\0A\00"
@kernel_0_name = private unnamed_addr constant [9 x i8] c"kernel_0\00"

; Function Attrs: sspstrong
define void @julia_copy_mat_64610(i8** dereferenceable(40), i8** dereferenceable(40)) #0 !dbg !5 {
top:
  %polly_launch_0_params = alloca [5 x i8*]
  %polly_launch_0_param_0 = alloca i8*
  %polly_launch_0_param_1 = alloca i8*
  %polly_launch_0_param_2 = alloca i64
  %polly_launch_0_param_3 = alloca i64
  %polly_launch_0_param_4 = alloca i64
  %polly_launch_0_params_i8ptr = bitcast [5 x i8*]* %polly_launch_0_params to i8*
  br label %top.split, !dbg !22

top.split:                                        ; preds = %top
  call void @llvm.dbg.value(metadata i8** null, i64 0, metadata !15, metadata !23), !dbg !22
  call void @llvm.dbg.value(metadata i8** null, i64 0, metadata !16, metadata !23), !dbg !22
  call void @llvm.dbg.value(metadata i8** %0, i64 0, metadata !15, metadata !23), !dbg !22
  %2 = bitcast i8** %0 to double**
  %3 = load double*, double** %2, align 8, !tbaa !24
  %4 = getelementptr i8*, i8** %0, i64 3
  %5 = bitcast i8** %4 to i64*
  %6 = load i64, i64* %5, align 8, !tbaa !24
  call void @llvm.dbg.value(metadata i8** %1, i64 0, metadata !16, metadata !23), !dbg !22
  %7 = bitcast i8** %1 to double**
  %8 = load double*, double** %7, align 8, !tbaa !24
  %9 = getelementptr i8*, i8** %1, i64 3
  %10 = bitcast i8** %9 to i64*
  %11 = load i64, i64* %10, align 8, !tbaa !24
  %12 = getelementptr i8*, i8** %0, i64 4, !dbg !22
  %13 = bitcast i8** %12 to i64*, !dbg !22
  %14 = load i64, i64* %13, align 8, !dbg !22, !tbaa !24
  call void @llvm.dbg.value(metadata i64 %6, i64 0, metadata !20, metadata !23), !dbg !22
  call void @llvm.dbg.value(metadata i64 %6, i64 0, metadata !20, metadata !23), !dbg !22
  %15 = icmp sgt i64 %6, 0, !dbg !27
  %16 = select i1 %15, i64 %6, i64 0, !dbg !27
  %17 = icmp slt i64 %6, 1, !dbg !27
  br i1 %17, label %L33, label %if.lr.ph, !dbg !27

if.lr.ph:                                         ; preds = %top.split
  br label %polly.split_new_and_old, !dbg !27

L20.L11.loopexit_crit_edge:                       ; preds = %if3
  br label %L11.loopexit, !dbg !28

L11.loopexit:                                     ; preds = %L20.L11.loopexit_crit_edge, %if
  %18 = icmp eq i64 %"#temp#1.08", %16, !dbg !27
  br i1 %18, label %L11.L33_crit_edge, label %if, !dbg !27

polly.split_new_and_old:                          ; preds = %if.lr.ph
  %19 = icmp sge i64 %11, %6
  %20 = icmp sle i64 %6, 0
  %21 = sext i1 %20 to i64
  %22 = icmp eq i64 0, %21
  %23 = and i1 %19, %22
  %polly.access.cast. = bitcast double* %8 to i64*
  %24 = call { i64, i1 } @llvm.ssub.with.overflow.i64(i64 %14, i64 1)
  %.obit = extractvalue { i64, i1 } %24, 1
  %polly.overflow.state = or i1 false, %.obit
  %.res = extractvalue { i64, i1 } %24, 0
  %polly.access.mul. = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 %.res, i64 %11)
  %polly.access.mul..obit = extractvalue { i64, i1 } %polly.access.mul., 1
  %polly.overflow.state9 = or i1 %polly.overflow.state, %polly.access.mul..obit
  %polly.access.mul..res = extractvalue { i64, i1 } %polly.access.mul., 0
  %polly.access.add. = call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %polly.access.mul..res, i64 %6)
  %polly.access.add..obit = extractvalue { i64, i1 } %polly.access.add., 1
  %polly.overflow.state10 = or i1 %polly.overflow.state9, %polly.access.add..obit
  %polly.access.add..res = extractvalue { i64, i1 } %polly.access.add., 0
  %polly.access. = getelementptr i64, i64* %polly.access.cast., i64 %polly.access.add..res
  %polly.access.cast.11 = bitcast double* %3 to i64*
  %polly.access.mul.12 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 0, i64 %6)
  %polly.access.mul..obit13 = extractvalue { i64, i1 } %polly.access.mul.12, 1
  %polly.overflow.state14 = or i1 %polly.overflow.state10, %polly.access.mul..obit13
  %polly.access.mul..res15 = extractvalue { i64, i1 } %polly.access.mul.12, 0
  %polly.access.add.16 = call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %polly.access.mul..res15, i64 0)
  %polly.access.add..obit17 = extractvalue { i64, i1 } %polly.access.add.16, 1
  %polly.overflow.state18 = or i1 %polly.overflow.state14, %polly.access.add..obit17
  %polly.access.add..res19 = extractvalue { i64, i1 } %polly.access.add.16, 0
  %polly.access.20 = getelementptr i64, i64* %polly.access.cast.11, i64 %polly.access.add..res19
  %25 = ptrtoint i64* %polly.access. to i64
  %26 = ptrtoint i64* %polly.access.20 to i64
  %27 = icmp ule i64 %25, %26
  %polly.access.cast.21 = bitcast double* %3 to i64*
  %28 = call { i64, i1 } @llvm.ssub.with.overflow.i64(i64 %14, i64 1)
  %.obit22 = extractvalue { i64, i1 } %28, 1
  %polly.overflow.state23 = or i1 %polly.overflow.state18, %.obit22
  %.res24 = extractvalue { i64, i1 } %28, 0
  %polly.access.mul.25 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 %.res24, i64 %6)
  %polly.access.mul..obit26 = extractvalue { i64, i1 } %polly.access.mul.25, 1
  %polly.overflow.state27 = or i1 %polly.overflow.state23, %polly.access.mul..obit26
  %polly.access.mul..res28 = extractvalue { i64, i1 } %polly.access.mul.25, 0
  %polly.access.add.29 = call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %polly.access.mul..res28, i64 %6)
  %polly.access.add..obit30 = extractvalue { i64, i1 } %polly.access.add.29, 1
  %polly.overflow.state31 = or i1 %polly.overflow.state27, %polly.access.add..obit30
  %polly.access.add..res32 = extractvalue { i64, i1 } %polly.access.add.29, 0
  %polly.access.33 = getelementptr i64, i64* %polly.access.cast.21, i64 %polly.access.add..res32
  %polly.access.cast.34 = bitcast double* %8 to i64*
  %polly.access.mul.35 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 0, i64 %11)
  %polly.access.mul..obit36 = extractvalue { i64, i1 } %polly.access.mul.35, 1
  %polly.overflow.state37 = or i1 %polly.overflow.state31, %polly.access.mul..obit36
  %polly.access.mul..res38 = extractvalue { i64, i1 } %polly.access.mul.35, 0
  %polly.access.add.39 = call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %polly.access.mul..res38, i64 0)
  %polly.access.add..obit40 = extractvalue { i64, i1 } %polly.access.add.39, 1
  %polly.overflow.state41 = or i1 %polly.overflow.state37, %polly.access.add..obit40
  %polly.access.add..res42 = extractvalue { i64, i1 } %polly.access.add.39, 0
  %polly.access.43 = getelementptr i64, i64* %polly.access.cast.34, i64 %polly.access.add..res42
  %29 = ptrtoint i64* %polly.access.33 to i64
  %30 = ptrtoint i64* %polly.access.43 to i64
  %31 = icmp ule i64 %29, %30
  %32 = or i1 %27, %31
  %33 = and i1 %23, %32
  %34 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 1, i64 %6)
  %.obit44 = extractvalue { i64, i1 } %34, 1
  %polly.overflow.state45 = or i1 %polly.overflow.state41, %.obit44
  %.res46 = extractvalue { i64, i1 } %34, 0
  %35 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 %.res46, i64 %14)
  %.obit47 = extractvalue { i64, i1 } %35, 1
  %polly.overflow.state48 = or i1 %polly.overflow.state45, %.obit47
  %.res49 = extractvalue { i64, i1 } %35, 0
  %36 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 20, i64 %.res49)
  %.obit50 = extractvalue { i64, i1 } %36, 1
  %polly.overflow.state51 = or i1 %polly.overflow.state48, %.obit50
  %.res52 = extractvalue { i64, i1 } %36, 0
  %37 = call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 0, i64 %.res52)
  %.obit53 = extractvalue { i64, i1 } %37, 1
  %polly.overflow.state54 = or i1 %polly.overflow.state51, %.obit53
  %.res55 = extractvalue { i64, i1 } %37, 0
  %38 = icmp sge i64 %.res55, 2621440
  %39 = and i1 %33, %38
  %polly.rtc.overflown = xor i1 %polly.overflow.state54, true
  %polly.rtc.result = and i1 %39, %polly.rtc.overflown
  br i1 %polly.rtc.result, label %polly.start, label %if.pre_entry_bb

if.pre_entry_bb:                                  ; preds = %polly.split_new_and_old
  br label %if, !dbg !27

if:                                               ; preds = %if.pre_entry_bb, %L11.loopexit
  %"#temp#1.08" = phi i64 [ %40, %L11.loopexit ], [ 1, %if.pre_entry_bb ]
  %40 = add i64 %"#temp#1.08", 1, !dbg !27
  call void @llvm.dbg.value(metadata i64 %14, i64 0, metadata !21, metadata !23), !dbg !22
  call void @llvm.dbg.value(metadata i64 %14, i64 0, metadata !21, metadata !23), !dbg !22
  %41 = icmp sgt i64 %14, 0, !dbg !28
  %42 = select i1 %41, i64 %14, i64 0, !dbg !28
  %43 = icmp slt i64 %14, 1, !dbg !28
  br i1 %43, label %L11.loopexit, label %if3.lr.ph, !dbg !28

if3.lr.ph:                                        ; preds = %if
  br label %if3, !dbg !28

L11.L33_crit_edge:                                ; preds = %L11.loopexit
  br label %polly.merge_new_and_old, !dbg !27

polly.merge_new_and_old:                          ; preds = %polly.exiting, %L11.L33_crit_edge
  br label %L33, !dbg !29

L33:                                              ; preds = %polly.merge_new_and_old, %top.split
  ret void, !dbg !29

if3:                                              ; preds = %if3.lr.ph, %if3
  %"#temp#.07" = phi i64 [ 1, %if3.lr.ph ], [ %44, %if3 ]
  %44 = add i64 %"#temp#.07", 1, !dbg !28
  call void @llvm.dbg.value(metadata i64 %"#temp#1.08", i64 0, metadata !19, metadata !23), !dbg !22
  call void @llvm.dbg.value(metadata i64 %"#temp#.07", i64 0, metadata !17, metadata !23), !dbg !22
  %45 = add i64 %"#temp#1.08", -1, !dbg !29
  %46 = add i64 %"#temp#.07", -1, !dbg !29
  %47 = mul i64 %46, %11, !dbg !29
  %48 = add i64 %45, %47, !dbg !29
  %49 = getelementptr double, double* %8, i64 %48, !dbg !29
  %50 = bitcast double* %49 to i64*, !dbg !29
  %51 = load i64, i64* %50, align 8, !dbg !29, !tbaa !30
  call void @llvm.dbg.value(metadata i64 %"#temp#1.08", i64 0, metadata !19, metadata !23), !dbg !22
  call void @llvm.dbg.value(metadata i64 %"#temp#.07", i64 0, metadata !17, metadata !23), !dbg !22
  %52 = mul i64 %46, %6, !dbg !29
  %53 = add i64 %45, %52, !dbg !29
  %54 = getelementptr double, double* %3, i64 %53, !dbg !29
  %55 = bitcast double* %54 to i64*, !dbg !29
  store i64 %51, i64* %55, align 8, !dbg !29, !tbaa !30
  %56 = icmp eq i64 %"#temp#.07", %42, !dbg !28
  br i1 %56, label %L20.L11.loopexit_crit_edge, label %if3, !dbg !28

polly.start:                                      ; preds = %polly.split_new_and_old
  br label %polly.acc.initialize

polly.acc.initialize:                             ; preds = %polly.start
  %57 = call i8* @polly_initContext()
  %58 = mul nsw i64 %14, %11
  %59 = mul i64 8, %58
  %p_dev_array_MemRef0 = call i8* @polly_allocateMemoryForDevice(i64 %59)
  %60 = mul nsw i64 %14, %6
  %61 = mul i64 8, %60
  %p_dev_array_MemRef1 = call i8* @polly_allocateMemoryForDevice(i64 %61)
  br label %polly.cond

polly.cond:                                       ; preds = %polly.acc.initialize
  %62 = icmp sge i64 %14, 1
  %63 = icmp sge i64 %6, 1
  %64 = and i1 %62, %63
  br i1 %64, label %polly.then, label %polly.else

polly.merge:                                      ; preds = %polly.else, %polly.merge57
  call void @polly_freeDeviceMemory(i8* %p_dev_array_MemRef0)
  call void @polly_freeDeviceMemory(i8* %p_dev_array_MemRef1)
  call void @polly_freeContext(i8* %57)
  br label %polly.exiting

polly.exiting:                                    ; preds = %polly.merge
  br label %polly.merge_new_and_old

polly.then:                                       ; preds = %polly.cond
  br label %polly.cond56

polly.cond56:                                     ; preds = %polly.then
  %65 = icmp sge i64 %11, 1
  br i1 %65, label %polly.then58, label %polly.else59

polly.merge57:                                    ; preds = %polly.else59, %polly.then58
  %66 = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef0)
  %67 = getelementptr [5 x i8*], [5 x i8*]* %polly_launch_0_params, i64 0, i64 0
  store i8* %66, i8** %polly_launch_0_param_0
  %68 = bitcast i8** %polly_launch_0_param_0 to i8*
  store i8* %68, i8** %67
  %69 = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef1)
  %70 = getelementptr [5 x i8*], [5 x i8*]* %polly_launch_0_params, i64 0, i64 1
  store i8* %69, i8** %polly_launch_0_param_1
  %71 = bitcast i8** %polly_launch_0_param_1 to i8*
  store i8* %71, i8** %70
  store i64 %14, i64* %polly_launch_0_param_2
  %72 = getelementptr [5 x i8*], [5 x i8*]* %polly_launch_0_params, i64 0, i64 2
  %73 = bitcast i64* %polly_launch_0_param_2 to i8*
  store i8* %73, i8** %72
  store i64 %6, i64* %polly_launch_0_param_3
  %74 = getelementptr [5 x i8*], [5 x i8*]* %polly_launch_0_params, i64 0, i64 3
  %75 = bitcast i64* %polly_launch_0_param_3 to i8*
  store i8* %75, i8** %74
  store i64 %11, i64* %polly_launch_0_param_4
  %76 = getelementptr [5 x i8*], [5 x i8*]* %polly_launch_0_params, i64 0, i64 4
  %77 = bitcast i64* %polly_launch_0_param_4 to i8*
  store i8* %77, i8** %76
  %78 = call i8* @polly_getKernel(i8* getelementptr inbounds ([2653 x i8], [2653 x i8]* @kernel_0, i32 0, i32 0), i8* getelementptr inbounds ([9 x i8], [9 x i8]* @kernel_0_name, i32 0, i32 0))
  %79 = icmp sge i64 %6, 8161
  %80 = add nsw i64 %6, 31
  %polly.fdiv_q.shr = ashr i64 %80, 5
  %81 = select i1 %79, i64 256, i64 %polly.fdiv_q.shr
  %82 = trunc i64 %81 to i32
  %83 = icmp sge i64 %14, 8162
  %84 = add nsw i64 %14, 31
  %polly.fdiv_q.shr60 = ashr i64 %84, 5
  %85 = select i1 %83, i64 256, i64 %polly.fdiv_q.shr60
  %86 = trunc i64 %85 to i32
  call void @polly_launchKernel(i8* %78, i32 %82, i32 %86, i32 32, i32 16, i32 1, i8* %polly_launch_0_params_i8ptr)
  call void @polly_freeKernel(i8* %78)
  %87 = mul nsw i64 %14, %6
  %88 = mul i64 8, %87
  %89 = bitcast double* %3 to i8*
  call void @polly_copyFromDeviceToHost(i8* %p_dev_array_MemRef1, i8* %89, i64 %88)
  br label %polly.merge

polly.else:                                       ; preds = %polly.cond
  br label %polly.merge

polly.then58:                                     ; preds = %polly.cond56
  %90 = mul nsw i64 %14, %11
  %91 = mul i64 8, %90
  %92 = bitcast double* %8 to i8*
  call void @polly_copyFromHostToDevice(i8* %92, i8* %p_dev_array_MemRef0, i64 %91)
  br label %polly.merge57

polly.else59:                                     ; preds = %polly.cond56
  br label %polly.merge57
}

define i8** @jlcall_copy_mat_64609(i8**, i8***, i32) #1 {
top:
  br label %top.split

top.split:                                        ; preds = %top
  %3 = load i8**, i8*** %1, align 8
  %4 = getelementptr i8**, i8*** %1, i64 1
  %5 = load i8**, i8*** %4, align 8
  call void @julia_copy_mat_64610(i8** %3, i8** %5)
  ret i8** inttoptr (i64 139756920668176 to i8**)
}

declare i8**** @jl_get_ptls_states()

; Function Attrs: nounwind readnone
declare void @llvm.dbg.declare(metadata, metadata, metadata) #2

; Function Attrs: argmemonly nounwind
declare void @llvm.memcpy.p0i8.p0i8.i32(i8* nocapture writeonly, i8* nocapture readonly, i32, i32, i1) #3

; Function Attrs: argmemonly nounwind
declare void @llvm.memcpy.p0i8.p0i8.i64(i8* nocapture writeonly, i8* nocapture readonly, i64, i32, i1) #3

declare void @jl_enter_handler(i8*)

; Function Attrs: returns_twice
declare i32 @__sigsetjmp(i8*, i32) #4

; Function Attrs: nounwind readnone
declare void @llvm.dbg.value(metadata, i64, metadata, metadata) #2

; Function Attrs: nounwind readnone
declare { i64, i1 } @llvm.ssub.with.overflow.i64(i64, i64) #2

; Function Attrs: nounwind readnone
declare { i64, i1 } @llvm.smul.with.overflow.i64(i64, i64) #2

; Function Attrs: nounwind readnone
declare { i64, i1 } @llvm.sadd.with.overflow.i64(i64, i64) #2

declare i8* @polly_initContext()

declare i8* @polly_allocateMemoryForDevice(i64)

declare void @polly_copyFromHostToDevice(i8*, i8*, i64)

declare i8* @polly_getDevicePtr(i8*)

declare i8* @polly_getKernel(i8*, i8*)

declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)

declare void @polly_freeKernel(i8*)

declare void @polly_copyFromDeviceToHost(i8*, i8*, i64)

declare void @polly_freeDeviceMemory(i8*)

declare void @polly_freeContext(i8*)

attributes #0 = { sspstrong "no-frame-pointer-elim"="true" }
attributes #1 = { "no-frame-pointer-elim"="true" }
attributes #2 = { nounwind readnone }
attributes #3 = { argmemonly nounwind }
attributes #4 = { returns_twice }

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

!0 = !{i32 2, !"Dwarf Version", i32 4}
!1 = !{i32 1, !"Debug Info Version", i32 3}
!2 = distinct !DICompileUnit(language: DW_LANG_C89, file: !3, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug, enums: !4)
!3 = !DIFile(filename: "REPL[19]", directory: ".")
!4 = !{}
!5 = distinct !DISubprogram(name: "copy_mat", linkageName: "julia_copy_mat_64610", scope: null, file: !3, type: !6, isLocal: false, isDefinition: true, isOptimized: true, unit: !2, variables: !12)
!6 = !DISubroutineType(types: !7)
!7 = !{!8, !8}
!8 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !9, size: 64, align: 64)
!9 = !DICompositeType(tag: DW_TAG_structure_type, name: "jl_value_t", file: !10, line: 71, align: 64, elements: !11)
!10 = !DIFile(filename: "julia.h", directory: "")
!11 = !{!8}
!12 = !{!13, !15, !16, !17, !19, !20, !21}
!13 = !DILocalVariable(name: "#self#", arg: 1, scope: !5, file: !3, line: 2, type: !14)
!14 = !DICompositeType(tag: DW_TAG_structure_type, name: "#copy_mat", align: 8, elements: !4, runtimeLang: DW_LANG_Julia, identifier: "#copy_mat_64588")
!15 = !DILocalVariable(name: "a", arg: 2, scope: !5, file: !3, line: 2, type: !8)
!16 = !DILocalVariable(name: "b", arg: 3, scope: !5, file: !3, line: 2, type: !8)
!17 = !DILocalVariable(name: "j", scope: !5, file: !3, line: 2, type: !18)
!18 = !DIBasicType(name: "Int64", size: 64, encoding: DW_ATE_unsigned)
!19 = !DILocalVariable(name: "i", scope: !5, file: !3, line: 2, type: !18)
!20 = !DILocalVariable(name: "ni", scope: !5, file: !3, line: 2, type: !18)
!21 = !DILocalVariable(name: "nj", scope: !5, file: !3, line: 2, type: !18)
!22 = !DILocation(line: 2, scope: !5)
!23 = !DIExpression()
!24 = !{!25, !25, i64 0, i64 1}
!25 = !{!"jtbaa_const", !26, i64 0}
!26 = !{!"jtbaa"}
!27 = !DILocation(line: 3, scope: !5)
!28 = !DILocation(line: 4, scope: !5)
!29 = !DILocation(line: 5, scope: !5)
!30 = !{!31, !31, i64 0}
!31 = !{!"jtbaa_arraybuf", !32, i64 0}
!32 = !{!"jtbaa_data", !26, i64 0}