Enzyme.jl
Enzyme.jl copied to clipboard
EnzymeAD
WIP: Fix GPUCompiler usage post v1.3
Ok this improves the situation for #2395 but is not yet enough
warning: linking module flags 'Dwarf Version': IDs have conflicting values ('i32 4' from start with 'i32 2' from start)
ERROR: LoadError: Failed to compile PTX code (ptxas exited with code 255)
Invocation arguments: --generate-line-info --verbose --gpu-name sm_75 --output-file /tmp/jl_sjSxgtgdgb.cubin /tmp/jl_CtIhwfKLEX.ptx
ptxas /tmp/jl_CtIhwfKLEX.ptx, line 959; error : Call has wrong number of parameters
ptxas /tmp/jl_CtIhwfKLEX.ptx, line 959; error : Type of argument does not match formal parameter 'julia_f_18874_param_1'
ptxas fatal : Ptx assembly aborted due to errors
If you think this is a bug, please file an issue and attach /tmp/jl_CtIhwfKLEX.ptx
//
// Generated by LLVM NVPTX Back-End
//
.version 8.8
.target sm_75
.address_size 64
.extern .func (.param .b32 func_retval0) vprintf
(
.param .b64 vprintf_param_0,
.param .b64 vprintf_param_1
)
;
.func gpu_report_exception
(
.param .align 16 .b8 gpu_report_exception_param_0[16],
.param .b64 gpu_report_exception_param_1
)
;
.func julia_arrayref_18878
(
.param .align 16 .b8 julia_arrayref_18878_param_0[16]
)
;
.func gpu_signal_exception
(
.param .align 16 .b8 gpu_signal_exception_param_0[16]
)
;
.global .align 1 .b8 __unnamed_1[107] = {83, 116, 97, 99, 107, 116, 114, 97, 99, 101, 32, 110, 111, 116, 32, 97, 118, 97, 105, 108, 97, 98, 108, 101, 44, 32, 114, 117, 110, 32, 74, 117, 108, 105, 97, 32, 111, 110, 32, 100, 101, 98, 117, 103, 32, 108, 101, 118, 101, 108, 32, 50, 32, 102, 111, 114, 32, 109, 111, 114, 101, 32, 100, 101, 116, 97, 105, 108, 115, 32, 40, 98, 121, 32, 112, 97, 115, 115, 105, 110, 103, 32, 45, 103, 50, 32, 116, 111, 32, 116, 104, 101, 32, 101, 120, 101, 99, 117, 116, 97, 98, 108, 101, 41, 46, 10, 0};
.global .align 1 .b8 __unnamed_2[4] = {37, 115, 10, 0};
.global .align 1 .b8 __unnamed_3[94] = {69, 82, 82, 79, 82, 58, 32, 97, 32, 37, 115, 32, 119, 97, 115, 32, 116, 104, 114, 111, 119, 110, 32, 100, 117, 114, 105, 110, 103, 32, 107, 101, 114, 110, 101, 108, 32, 101, 120, 101, 99, 117, 116, 105, 111, 110, 32, 111, 110, 32, 116, 104, 114, 101, 97, 100, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 32, 105, 110, 32, 98, 108, 111, 99, 107, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 46, 10, 0};
.global .align 1 .b8 __unnamed_4[2] = {10, 0};
.global .align 1 .b8 exception16[10] = {101, 120, 99, 101, 112, 116, 105, 111, 110, 0};
// -- Begin function julia_error_18883
// @julia_error_18883
.func julia_error_18883(
.param .align 16 .b8 julia_error_18883_param_0[16]
)
{
.reg .b32 %r<2>;
.reg .b64 %rd<4>;
.loc 1 35 0 // error.jl:35:0
$L__func_begin0:
.loc 1 35 0 // error.jl:35:0
// %bb.0: // %top
ld.param.u64 %rd1, [julia_error_18883_param_0];
ld.param.u32 %r1, [julia_error_18883_param_0+8];
$L__tmp0:
.loc 1 35 0 // error.jl:35:0
mov.u64 %rd2, exception16;
cvta.global.u64 %rd3, %rd2;
{ // callseq 8, 0
.reg .b32 temp_param_reg;
.param .align 16 .b8 param0[16];
st.param.b64 [param0+0], %rd1;
st.param.b32 [param0+8], %r1;
.param .b64 param1;
st.param.b64 [param1+0], %rd3;
call.uni
gpu_report_exception,
(
param0,
param1
);
} // callseq 8
{ // callseq 9, 0
.reg .b32 temp_param_reg;
.param .align 16 .b8 param0[16];
st.param.b64 [param0+0], %rd1;
st.param.b32 [param0+8], %r1;
call.uni
gpu_signal_exception,
(
param0
);
} // callseq 9
trap;
trap;
// begin inline asm
exit;
// end inline asm
$L__tmp1:
$L__func_end0:
// -- End function
}
.func gpu_report_exception(
.param .align 16 .b8 gpu_report_exception_param_0[16],
.param .b64 gpu_report_exception_param_1
) // -- Begin function gpu_report_exception
// @gpu_report_exception
{
.local .align 8 .b8 __local_depot1[32];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<14>;
.reg .b32 %r<126>;
.reg .b64 %rd<67>;
.loc 2 143 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:143:0
$L__func_begin1:
.loc 2 143 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:143:0
// %bb.0: // %top
mov.u64 %SPL, __local_depot1;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd4, [gpu_report_exception_param_0];
$L__tmp2:
.loc 3 316 0 // pointer.jl:316:0
add.s64 %rd11, %rd4, 4;
$L__tmp3:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
atom.cas.b32 %r19, [%rd11], 0, 1;
$L__tmp4:
.loc 2 129 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:129:0
setp.ne.s32 %p1, %r19, 0;
@%p1 bra $L__BB1_2;
// %bb.1: // %L12
$L__tmp5:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r96, %tid.x;
$L__tmp6:
.loc 5 87 0 // int.jl:87:0
add.s32 %r125, %r96, 1;
$L__tmp7:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r97, %tid.y;
$L__tmp8:
.loc 5 87 0 // int.jl:87:0
add.s32 %r124, %r97, 1;
$L__tmp9:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r98, %tid.z;
$L__tmp10:
.loc 5 87 0 // int.jl:87:0
add.s32 %r123, %r98, 1;
$L__tmp11:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r99, %ctaid.x;
$L__tmp12:
.loc 5 87 0 // int.jl:87:0
add.s32 %r122, %r99, 1;
$L__tmp13:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r100, %ctaid.y;
$L__tmp14:
.loc 5 87 0 // int.jl:87:0
add.s32 %r121, %r100, 1;
$L__tmp15:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r101, %ctaid.z;
$L__tmp16:
.loc 5 87 0 // int.jl:87:0
add.s32 %r120, %r101, 1;
mov.u32 %r102, 0;
$L__tmp17:
.loc 3 180 0 // pointer.jl:180:0
st.u8 [%rd4+11], %r102;
st.u8 [%rd4+10], %r102;
st.u8 [%rd4+8], %r125;
shr.u32 %r103, %r125, 8;
st.u8 [%rd4+9], %r103;
st.u8 [%rd4+15], %r102;
st.u8 [%rd4+14], %r102;
st.u8 [%rd4+12], %r124;
shr.u32 %r104, %r124, 8;
st.u8 [%rd4+13], %r104;
st.u8 [%rd4+19], %r102;
st.u8 [%rd4+18], %r102;
st.u8 [%rd4+17], %r102;
st.u8 [%rd4+16], %r123;
$L__tmp18:
.loc 3 180 0 // pointer.jl:180:0
st.u8 [%rd4+20], %r122;
shr.u32 %r105, %r122, 24;
st.u8 [%rd4+23], %r105;
shr.u32 %r106, %r122, 16;
st.u8 [%rd4+22], %r106;
shr.u32 %r107, %r122, 8;
st.u8 [%rd4+21], %r107;
st.u8 [%rd4+27], %r102;
st.u8 [%rd4+24], %r121;
shr.u32 %r108, %r121, 16;
st.u8 [%rd4+26], %r108;
shr.u32 %r109, %r121, 8;
st.u8 [%rd4+25], %r109;
st.u8 [%rd4+31], %r102;
st.u8 [%rd4+28], %r120;
shr.u32 %r110, %r120, 16;
st.u8 [%rd4+30], %r110;
shr.u32 %r111, %r120, 8;
st.u8 [%rd4+29], %r111;
$L__tmp19:
.loc 6 98 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/synchronization.jl:98:0
membar.gl;
$L__tmp20:
$L__BB1_7: // %L132.critedge
.loc 6 0 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/synchronization.jl:0:0
ld.param.u64 %rd7, [gpu_report_exception_param_1];
add.u64 %rd8, %SP, 0;
add.u64 %rd1, %SPL, 0;
$L__tmp21:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %rd12, [%rd4+32];
ld.u8 %rd13, [%rd4+33];
shl.b64 %rd14, %rd13, 8;
or.b64 %rd15, %rd14, %rd12;
ld.u8 %rd16, [%rd4+34];
shl.b64 %rd17, %rd16, 16;
ld.u8 %rd18, [%rd4+35];
shl.b64 %rd19, %rd18, 24;
or.b64 %rd20, %rd19, %rd17;
or.b64 %rd21, %rd20, %rd15;
ld.u8 %rd22, [%rd4+36];
ld.u8 %rd23, [%rd4+37];
shl.b64 %rd24, %rd23, 8;
or.b64 %rd25, %rd24, %rd22;
ld.u8 %rd26, [%rd4+38];
shl.b64 %rd27, %rd26, 16;
ld.u8 %rd28, [%rd4+39];
shl.b64 %rd29, %rd28, 24;
or.b64 %rd30, %rd29, %rd27;
or.b64 %rd31, %rd30, %rd25;
shl.b64 %rd32, %rd31, 32;
or.b64 %rd33, %rd32, %rd21;
$L__tmp22:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s64 %p11, %rd33, 0;
setp.eq.s64 %p12, %rd33, 0;
$L__tmp23:
.loc 2 148 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:148:0
selp.b64 %rd5, %rd7, %rd33, %p12;
.loc 2 151 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:151:0
@%p11 bra $L__BB1_9;
bra.uni $L__BB1_8;
$L__BB1_9: // %L204.cont
$L__tmp24:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
st.local.u64 [%rd1], %rd5;
st.local.v2.u32 [%rd1+8], {%r125, %r124};
st.local.v2.u32 [%rd1+16], {%r123, %r122};
st.local.v2.u32 [%rd1+24], {%r121, %r120};
mov.u64 %rd34, __unnamed_3;
cvta.global.u64 %rd35, %rd34;
{ // callseq 10, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd35;
.param .b64 param1;
st.param.b64 [param1+0], %rd8;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r112, [retval0+0];
} // callseq 10
bra.uni $L__BB1_10;
$L__tmp25:
$L__BB1_2: // %L44
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r20, [%rd4+4];
ld.u8 %r21, [%rd4+5];
shl.b32 %r22, %r21, 8;
or.b32 %r23, %r22, %r20;
ld.u8 %r24, [%rd4+6];
shl.b32 %r25, %r24, 16;
ld.u8 %r26, [%rd4+7];
shl.b32 %r27, %r26, 24;
or.b32 %r28, %r27, %r25;
or.b32 %r29, %r28, %r23;
$L__tmp26:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p2, %r29, 1;
$L__tmp27:
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
@%p2 bra $L__BB1_13;
// %bb.3: // %L52
$L__tmp28:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r30, [%rd4+8];
ld.u8 %r31, [%rd4+9];
shl.b32 %r32, %r31, 8;
or.b32 %r33, %r32, %r30;
ld.u8 %r34, [%rd4+10];
shl.b32 %r35, %r34, 16;
ld.u8 %r36, [%rd4+11];
shl.b32 %r37, %r36, 24;
or.b32 %r38, %r37, %r35;
or.b32 %r125, %r38, %r33;
$L__tmp29:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r39, %tid.x;
$L__tmp30:
.loc 5 87 0 // int.jl:87:0
add.s32 %r40, %r39, 1;
$L__tmp31:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p3, %r125, %r40;
$L__tmp32:
.loc 8 552 0 // tuple.jl:552:0
@%p3 bra $L__BB1_13;
$L__tmp33:
// %bb.4: // %L73
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
mov.u32 %r41, %tid.z;
add.s32 %r123, %r41, 1;
$L__tmp34:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r42, [%rd4+16];
ld.u8 %r43, [%rd4+17];
shl.b32 %r44, %r43, 8;
or.b32 %r45, %r44, %r42;
ld.u8 %r46, [%rd4+18];
shl.b32 %r47, %r46, 16;
ld.u8 %r48, [%rd4+19];
shl.b32 %r49, %r48, 24;
or.b32 %r50, %r49, %r47;
or.b32 %r51, %r50, %r45;
$L__tmp35:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r52, %tid.y;
$L__tmp36:
.loc 5 87 0 // int.jl:87:0
add.s32 %r124, %r52, 1;
$L__tmp37:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r53, [%rd4+12];
ld.u8 %r54, [%rd4+13];
shl.b32 %r55, %r54, 8;
or.b32 %r56, %r55, %r53;
ld.u8 %r57, [%rd4+14];
shl.b32 %r58, %r57, 16;
ld.u8 %r59, [%rd4+15];
shl.b32 %r60, %r59, 24;
or.b32 %r61, %r60, %r58;
or.b32 %r62, %r61, %r56;
$L__tmp38:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p4, %r62, %r124;
setp.ne.s32 %p5, %r51, %r123;
$L__tmp39:
.loc 8 552 0 // tuple.jl:552:0
or.pred %p6, %p4, %p5;
$L__tmp40:
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
@%p6 bra $L__BB1_13;
// %bb.5: // %L90
$L__tmp41:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r63, [%rd4+20];
ld.u8 %r64, [%rd4+21];
shl.b32 %r65, %r64, 8;
or.b32 %r66, %r65, %r63;
ld.u8 %r67, [%rd4+22];
shl.b32 %r68, %r67, 16;
ld.u8 %r69, [%rd4+23];
shl.b32 %r70, %r69, 24;
or.b32 %r71, %r70, %r68;
or.b32 %r122, %r71, %r66;
$L__tmp42:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r72, %ctaid.x;
$L__tmp43:
.loc 5 87 0 // int.jl:87:0
add.s32 %r73, %r72, 1;
$L__tmp44:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p7, %r122, %r73;
$L__tmp45:
.loc 8 552 0 // tuple.jl:552:0
@%p7 bra $L__BB1_13;
$L__tmp46:
// %bb.6: // %L111
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r74, %ctaid.z;
$L__tmp47:
.loc 5 87 0 // int.jl:87:0
add.s32 %r120, %r74, 1;
$L__tmp48:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r75, [%rd4+28];
ld.u8 %r76, [%rd4+29];
shl.b32 %r77, %r76, 8;
or.b32 %r78, %r77, %r75;
ld.u8 %r79, [%rd4+30];
shl.b32 %r80, %r79, 16;
ld.u8 %r81, [%rd4+31];
shl.b32 %r82, %r81, 24;
or.b32 %r83, %r82, %r80;
or.b32 %r84, %r83, %r78;
$L__tmp49:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r85, %ctaid.y;
$L__tmp50:
.loc 5 87 0 // int.jl:87:0
add.s32 %r121, %r85, 1;
$L__tmp51:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r86, [%rd4+24];
ld.u8 %r87, [%rd4+25];
shl.b32 %r88, %r87, 8;
or.b32 %r89, %r88, %r86;
ld.u8 %r90, [%rd4+26];
shl.b32 %r91, %r90, 16;
ld.u8 %r92, [%rd4+27];
shl.b32 %r93, %r92, 24;
or.b32 %r94, %r93, %r91;
or.b32 %r95, %r94, %r89;
$L__tmp52:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p8, %r95, %r121;
setp.ne.s32 %p9, %r84, %r120;
$L__tmp53:
.loc 8 552 0 // tuple.jl:552:0
or.pred %p10, %p8, %p9;
$L__tmp54:
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
@%p10 bra $L__BB1_13;
bra.uni $L__BB1_7;
$L__tmp55:
$L__BB1_8: // %L200.cont
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
st.local.u64 [%rd1], %rd5;
st.local.v2.u32 [%rd1+8], {%r125, %r124};
st.local.v2.u32 [%rd1+16], {%r123, %r122};
st.local.v2.u32 [%rd1+24], {%r121, %r120};
mov.u64 %rd37, __unnamed_3;
cvta.global.u64 %rd38, %rd37;
{ // callseq 11, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd38;
.param .b64 param1;
st.param.b64 [param1+0], %rd8;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r114, [retval0+0];
} // callseq 11
$L__tmp56:
$L__BB1_10: // %L209
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %rd40, [%rd4+40];
ld.u8 %rd41, [%rd4+41];
shl.b64 %rd42, %rd41, 8;
or.b64 %rd43, %rd42, %rd40;
ld.u8 %rd44, [%rd4+42];
shl.b64 %rd45, %rd44, 16;
ld.u8 %rd46, [%rd4+43];
shl.b64 %rd47, %rd46, 24;
or.b64 %rd48, %rd47, %rd45;
or.b64 %rd49, %rd48, %rd43;
ld.u8 %rd50, [%rd4+44];
ld.u8 %rd51, [%rd4+45];
shl.b64 %rd52, %rd51, 8;
or.b64 %rd53, %rd52, %rd50;
ld.u8 %rd54, [%rd4+46];
shl.b64 %rd55, %rd54, 16;
ld.u8 %rd56, [%rd4+47];
shl.b64 %rd57, %rd56, 24;
or.b64 %rd58, %rd57, %rd55;
or.b64 %rd59, %rd58, %rd53;
shl.b64 %rd60, %rd59, 32;
or.b64 %rd6, %rd60, %rd49;
$L__tmp57:
.loc 7 639 0 // promotion.jl:639:0
setp.eq.s64 %p13, %rd6, 0;
$L__tmp58:
.loc 2 153 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:153:0
@%p13 bra $L__BB1_12;
// %bb.11: // %L218
$L__tmp59:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
st.local.u64 [%rd1], %rd6;
mov.u64 %rd61, __unnamed_2;
cvta.global.u64 %rd62, %rd61;
{ // callseq 12, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd62;
.param .b64 param1;
st.param.b64 [param1+0], %rd8;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r116, [retval0+0];
} // callseq 12
$L__tmp60:
$L__BB1_12: // %L225
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u64 %rd64, __unnamed_1;
cvta.global.u64 %rd65, %rd64;
mov.u64 %rd66, 0;
{ // callseq 13, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd65;
.param .b64 param1;
st.param.b64 [param1+0], %rd66;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r118, [retval0+0];
} // callseq 13
$L__tmp61:
$L__BB1_13: // %L227
.loc 2 158 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:158:0
ret;
$L__tmp62:
$L__func_end1:
// -- End function
}
// .globl julia_f_18874 // -- Begin function julia_f_18874
.visible .func julia_f_18874(
.param .align 8 .b8 julia_f_18874_param_0[16],
.param .b64 julia_f_18874_param_1
) // @julia_f_18874
{
.reg .b32 %r<2>;
.reg .b64 %rd<2>;
.loc 14 4 0 // /home/vchuravy/src/Enzyme/CUDA.jl:4:0
$L__func_begin2:
.loc 14 4 0 // /home/vchuravy/src/Enzyme/CUDA.jl:4:0
// %bb.0: // %top
ld.param.u64 %rd1, [julia_f_18874_param_0];
ld.param.u32 %r1, [julia_f_18874_param_0+8];
$L__tmp63:
.loc 15 175 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/array.jl:175:0
{ // callseq 14, 0
.reg .b32 temp_param_reg;
.param .align 16 .b8 param0[16];
st.param.b64 [param0+0], %rd1;
st.param.b32 [param0+8], %r1;
call.uni
julia_arrayref_18878,
(
param0
);
} // callseq 14
trap;
// begin inline asm
exit;
// end inline asm
$L__tmp64:
$L__func_end2:
// -- End function
}
.func julia_arrayref_18878(
.param .align 16 .b8 julia_arrayref_18878_param_0[16]
) // -- Begin function julia_arrayref_18878
// @julia_arrayref_18878
{
.reg .b32 %r<2>;
.reg .b64 %rd<2>;
.loc 16 33 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/utils.jl:33:0
$L__func_begin3:
.loc 16 33 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/utils.jl:33:0
// %bb.0: // %top
ld.param.u64 %rd1, [julia_arrayref_18878_param_0];
ld.param.u32 %r1, [julia_arrayref_18878_param_0+8];
$L__tmp65:
.loc 16 33 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/utils.jl:33:0
{ // callseq 15, 0
.reg .b32 temp_param_reg;
.param .align 16 .b8 param0[16];
st.param.b64 [param0+0], %rd1;
st.param.b32 [param0+8], %r1;
call.uni
julia_error_18883,
(
param0
);
} // callseq 15
trap;
// begin inline asm
exit;
// end inline asm
$L__tmp66:
$L__func_end3:
// -- End function
}
.func gpu_signal_exception(
.param .align 16 .b8 gpu_signal_exception_param_0[16]
) // -- Begin function gpu_signal_exception
// @gpu_signal_exception
{
.reg .pred %p<11>;
.reg .b32 %r<113>;
.reg .b64 %rd<6>;
.loc 2 189 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:189:0
$L__func_begin4:
.loc 2 189 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:189:0
// %bb.0: // %top
ld.param.u64 %rd1, [gpu_signal_exception_param_0];
$L__tmp67:
.loc 3 316 0 // pointer.jl:316:0
add.s64 %rd2, %rd1, 4;
$L__tmp68:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
atom.cas.b32 %r1, [%rd2], 0, 1;
$L__tmp69:
.loc 2 129 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:129:0
setp.ne.s32 %p1, %r1, 0;
mov.u32 %r112, 0;
@%p1 bra $L__BB4_2;
// %bb.1: // %L12
$L__tmp70:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r84, %tid.x;
$L__tmp71:
.loc 5 87 0 // int.jl:87:0
add.s32 %r85, %r84, 1;
$L__tmp72:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r86, %tid.y;
$L__tmp73:
.loc 5 87 0 // int.jl:87:0
add.s32 %r87, %r86, 1;
$L__tmp74:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r88, %tid.z;
$L__tmp75:
.loc 5 87 0 // int.jl:87:0
add.s32 %r89, %r88, 1;
$L__tmp76:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r90, %ctaid.x;
$L__tmp77:
.loc 5 87 0 // int.jl:87:0
add.s32 %r91, %r90, 1;
$L__tmp78:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r92, %ctaid.y;
$L__tmp79:
.loc 5 87 0 // int.jl:87:0
add.s32 %r93, %r92, 1;
$L__tmp80:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r94, %ctaid.z;
$L__tmp81:
.loc 5 87 0 // int.jl:87:0
add.s32 %r95, %r94, 1;
$L__tmp82:
.loc 3 180 0 // pointer.jl:180:0
st.u8 [%rd1+11], %r112;
st.u8 [%rd1+10], %r112;
st.u8 [%rd1+8], %r85;
shr.u32 %r97, %r85, 8;
st.u8 [%rd1+9], %r97;
st.u8 [%rd1+15], %r112;
st.u8 [%rd1+14], %r112;
st.u8 [%rd1+12], %r87;
shr.u32 %r98, %r87, 8;
st.u8 [%rd1+13], %r98;
st.u8 [%rd1+19], %r112;
st.u8 [%rd1+18], %r112;
st.u8 [%rd1+17], %r112;
st.u8 [%rd1+16], %r89;
$L__tmp83:
.loc 3 180 0 // pointer.jl:180:0
st.u8 [%rd1+20], %r91;
shr.u32 %r99, %r91, 24;
st.u8 [%rd1+23], %r99;
shr.u32 %r100, %r91, 16;
st.u8 [%rd1+22], %r100;
shr.u32 %r101, %r91, 8;
st.u8 [%rd1+21], %r101;
st.u8 [%rd1+27], %r112;
st.u8 [%rd1+24], %r93;
shr.u32 %r102, %r93, 16;
st.u8 [%rd1+26], %r102;
shr.u32 %r103, %r93, 8;
st.u8 [%rd1+25], %r103;
st.u8 [%rd1+31], %r112;
st.u8 [%rd1+28], %r95;
shr.u32 %r104, %r95, 16;
st.u8 [%rd1+30], %r104;
shr.u32 %r105, %r95, 8;
st.u8 [%rd1+29], %r105;
$L__tmp84:
.loc 6 98 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/synchronization.jl:98:0
membar.gl;
$L__tmp85:
$L__BB4_7: // %L132.critedge
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u64 %rd3, __unnamed_4;
cvta.global.u64 %rd4, %rd3;
mov.u64 %rd5, 0;
{ // callseq 16, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd4;
.param .b64 param1;
st.param.b64 [param1+0], %rd5;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r106, [retval0+0];
} // callseq 16
$L__tmp86:
.loc 3 180 0 // pointer.jl:180:0
st.u8 [%rd1+7], %r112;
st.u8 [%rd1+6], %r112;
st.u8 [%rd1+5], %r112;
mov.u32 %r109, 2;
st.u8 [%rd1+4], %r109;
bra.uni $L__BB4_8;
$L__tmp87:
$L__BB4_2: // %L44
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r2, [%rd1+4];
ld.u8 %r3, [%rd1+5];
shl.b32 %r4, %r3, 8;
or.b32 %r5, %r4, %r2;
ld.u8 %r6, [%rd1+6];
shl.b32 %r7, %r6, 16;
ld.u8 %r8, [%rd1+7];
shl.b32 %r9, %r8, 24;
or.b32 %r10, %r9, %r7;
or.b32 %r11, %r10, %r5;
$L__tmp88:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p2, %r11, 1;
$L__tmp89:
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
@%p2 bra $L__BB4_8;
// %bb.3: // %L52
$L__tmp90:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r12, [%rd1+8];
ld.u8 %r13, [%rd1+9];
shl.b32 %r14, %r13, 8;
or.b32 %r15, %r14, %r12;
ld.u8 %r16, [%rd1+10];
shl.b32 %r17, %r16, 16;
ld.u8 %r18, [%rd1+11];
shl.b32 %r19, %r18, 24;
or.b32 %r20, %r19, %r17;
or.b32 %r21, %r20, %r15;
$L__tmp91:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r22, %tid.x;
$L__tmp92:
.loc 5 87 0 // int.jl:87:0
add.s32 %r23, %r22, 1;
$L__tmp93:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p3, %r21, %r23;
$L__tmp94:
.loc 8 552 0 // tuple.jl:552:0
@%p3 bra $L__BB4_8;
$L__tmp95:
// %bb.4: // %L73
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r24, %tid.z;
$L__tmp96:
.loc 5 87 0 // int.jl:87:0
add.s32 %r25, %r24, 1;
$L__tmp97:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r26, [%rd1+16];
ld.u8 %r27, [%rd1+17];
shl.b32 %r28, %r27, 8;
or.b32 %r29, %r28, %r26;
ld.u8 %r30, [%rd1+18];
shl.b32 %r31, %r30, 16;
ld.u8 %r32, [%rd1+19];
shl.b32 %r33, %r32, 24;
or.b32 %r34, %r33, %r31;
or.b32 %r35, %r34, %r29;
$L__tmp98:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r36, %tid.y;
$L__tmp99:
.loc 5 87 0 // int.jl:87:0
add.s32 %r37, %r36, 1;
$L__tmp100:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r38, [%rd1+12];
ld.u8 %r39, [%rd1+13];
shl.b32 %r40, %r39, 8;
or.b32 %r41, %r40, %r38;
ld.u8 %r42, [%rd1+14];
shl.b32 %r43, %r42, 16;
ld.u8 %r44, [%rd1+15];
shl.b32 %r45, %r44, 24;
or.b32 %r46, %r45, %r43;
or.b32 %r47, %r46, %r41;
$L__tmp101:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p4, %r47, %r37;
setp.ne.s32 %p5, %r35, %r25;
$L__tmp102:
.loc 8 552 0 // tuple.jl:552:0
or.pred %p6, %p4, %p5;
$L__tmp103:
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
@%p6 bra $L__BB4_8;
// %bb.5: // %L90
$L__tmp104:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r48, [%rd1+20];
ld.u8 %r49, [%rd1+21];
shl.b32 %r50, %r49, 8;
or.b32 %r51, %r50, %r48;
ld.u8 %r52, [%rd1+22];
shl.b32 %r53, %r52, 16;
ld.u8 %r54, [%rd1+23];
shl.b32 %r55, %r54, 24;
or.b32 %r56, %r55, %r53;
or.b32 %r57, %r56, %r51;
$L__tmp105:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r58, %ctaid.x;
$L__tmp106:
.loc 5 87 0 // int.jl:87:0
add.s32 %r59, %r58, 1;
$L__tmp107:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p7, %r57, %r59;
$L__tmp108:
.loc 8 552 0 // tuple.jl:552:0
@%p7 bra $L__BB4_8;
$L__tmp109:
// %bb.6: // %L111
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r60, %ctaid.z;
$L__tmp110:
.loc 5 87 0 // int.jl:87:0
add.s32 %r61, %r60, 1;
$L__tmp111:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r62, [%rd1+28];
ld.u8 %r63, [%rd1+29];
shl.b32 %r64, %r63, 8;
or.b32 %r65, %r64, %r62;
ld.u8 %r66, [%rd1+30];
shl.b32 %r67, %r66, 16;
ld.u8 %r68, [%rd1+31];
shl.b32 %r69, %r68, 24;
or.b32 %r70, %r69, %r67;
or.b32 %r71, %r70, %r65;
$L__tmp112:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
mov.u32 %r72, %ctaid.y;
$L__tmp113:
.loc 5 87 0 // int.jl:87:0
add.s32 %r73, %r72, 1;
$L__tmp114:
.loc 3 153 0 // pointer.jl:153:0
ld.u8 %r74, [%rd1+24];
ld.u8 %r75, [%rd1+25];
shl.b32 %r76, %r75, 8;
or.b32 %r77, %r76, %r74;
ld.u8 %r78, [%rd1+26];
shl.b32 %r79, %r78, 16;
ld.u8 %r80, [%rd1+27];
shl.b32 %r81, %r80, 24;
or.b32 %r82, %r81, %r79;
or.b32 %r83, %r82, %r77;
$L__tmp115:
.loc 7 639 0 // promotion.jl:639:0
setp.ne.s32 %p8, %r83, %r73;
setp.ne.s32 %p9, %r71, %r61;
$L__tmp116:
.loc 8 552 0 // tuple.jl:552:0
or.pred %p10, %p8, %p9;
$L__tmp117:
.loc 2 134 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:134:0
@%p10 bra $L__BB4_8;
bra.uni $L__BB4_7;
$L__tmp118:
$L__BB4_8: // %L139
.loc 3 180 0 // pointer.jl:180:0
st.u8 [%rd1+3], %r112;
st.u8 [%rd1+2], %r112;
st.u8 [%rd1+1], %r112;
mov.u32 %r111, 1;
st.u8 [%rd1], %r111;
$L__tmp119:
.loc 6 109 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/synchronization.jl:109:0
membar.sys;
$L__tmp120:
.loc 4 39 0 // /home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl:39:0
// begin inline asm
exit;
// end inline asm
$L__tmp121:
.loc 2 205 0 // /home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl:205:0
ret;
$L__tmp122:
$L__func_end4:
// -- End function
}
// .globl _Z2df13CuDeviceArrayI7Float32Li1ELi1EES1_ // -- Begin function _Z2df13CuDeviceArrayI7Float32Li1ELi1EES1_
.visible .entry _Z2df13CuDeviceArrayI7Float32Li1ELi1EES1_(
.param .align 8 .b8 _Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_0[16],
.param .align 8 .b8 _Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_1[32],
.param .align 8 .b8 _Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_2[32]
) // @_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1_
{
.reg .b32 %r<2>;
.reg .b64 %rd<10>;
.loc 14 13 0 // /home/vchuravy/src/Enzyme/CUDA.jl:13:0
$L__func_begin5:
.loc 14 13 0 // /home/vchuravy/src/Enzyme/CUDA.jl:13:0
// %bb.0: // %conversion
ld.param.u64 %rd1, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_0];
ld.param.u32 %r1, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_0+8];
ld.param.u64 %rd2, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_1];
ld.param.u64 %rd3, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_1+8];
ld.param.u64 %rd4, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_1+16];
ld.param.u64 %rd5, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_1+24];
ld.param.u64 %rd6, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_2];
ld.param.u64 %rd7, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_2+8];
ld.param.u64 %rd8, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_2+16];
ld.param.u64 %rd9, [_Z2df13CuDeviceArrayI7Float32Li1ELi1EES1__param_2+24];
$L__tmp123:
.loc 19 5485 0 // /home/vchuravy/src/Enzyme/src/compiler.jl:5485:0
{ // callseq 17, 0
.reg .b32 temp_param_reg;
.param .align 8 .b8 param0[16];
st.param.b64 [param0+0], %rd1;
st.param.b32 [param0+8], %r1;
.param .align 8 .b8 param1[32];
st.param.b64 [param1+0], %rd2;
st.param.b64 [param1+8], %rd3;
st.param.b64 [param1+16], %rd4;
st.param.b64 [param1+24], %rd5;
.param .align 8 .b8 param2[1];
st.param.b64 [param2+0], %rd6;
st.param.b64 [param2+8], %rd7;
st.param.b64 [param2+16], %rd8;
st.param.b64 [param2+24], %rd9;
call.uni
julia_f_18874,
(
param0,
param1,
param2
);
} // callseq 17
$L__tmp124:
.loc 14 15 0 // /home/vchuravy/src/Enzyme/CUDA.jl:15:0
trap;
// begin inline asm
exit;
// end inline asm
$L__tmp125:
$L__func_end5:
// -- End function
}
.file 1 "./error.jl"
.file 2 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/runtime.jl"
.file 3 "./pointer.jl"
.file 4 "/home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/base.jl"
.file 5 "./int.jl"
.file 6 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/synchronization.jl"
.file 7 "./promotion.jl"
.file 8 "./tuple.jl"
.file 9 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/atomics.jl"
.file 10 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/indexing.jl"
.file 11 "./namedtuple.jl"
.file 12 "./operators.jl"
.file 13 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/output.jl"
.file 14 "/home/vchuravy/src/Enzyme/CUDA.jl"
.file 15 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/array.jl"
.file 16 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/utils.jl"
.file 17 "/home/vchuravy/.julia/packages/CUDA/LhtzZ/src/device/intrinsics/misc.jl"
.file 18 "/home/vchuravy/.julia/packages/LLVM/2JPxT/src/interop/asmcall.jl"
.file 19 "/home/vchuravy/src/Enzyme/src/compiler.jl"
.file 20 "/home/vchuravy/src/Enzyme/src/Enzyme.jl"
.section .debug_abbrev
{
.b8 1 // Abbreviation Code
.b8 17 // DW_TAG_compile_unit
.b8 0 // DW_CHILDREN_no
.b8 37 // DW_AT_producer
.b8 8 // DW_FORM_string
.b8 19 // DW_AT_language
.b8 5 // DW_FORM_data2
.b8 3 // DW_AT_name
.b8 8 // DW_FORM_string
.b8 16 // DW_AT_stmt_list
.b8 6 // DW_FORM_data4
.b8 27 // DW_AT_comp_dir
.b8 8 // DW_FORM_string
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 2 // Abbreviation Code
.b8 17 // DW_TAG_compile_unit
.b8 1 // DW_CHILDREN_yes
.b8 37 // DW_AT_producer
.b8 8 // DW_FORM_string
.b8 19 // DW_AT_language
.b8 5 // DW_FORM_data2
.b8 3 // DW_AT_name
.b8 8 // DW_FORM_string
.b8 16 // DW_AT_stmt_list
.b8 6 // DW_FORM_data4
.b8 27 // DW_AT_comp_dir
.b8 8 // DW_FORM_string
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 3 // Abbreviation Code
.b8 46 // DW_TAG_subprogram
.b8 0 // DW_CHILDREN_no
.b8 3 // DW_AT_name
.b8 8 // DW_FORM_string
.b8 32 // DW_AT_inline
.b8 11 // DW_FORM_data1
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 4 // Abbreviation Code
.b8 46 // DW_TAG_subprogram
.b8 1 // DW_CHILDREN_yes
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 3 // DW_AT_name
.b8 8 // DW_FORM_string
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 5 // Abbreviation Code
.b8 29 // DW_TAG_inlined_subroutine
.b8 1 // DW_CHILDREN_yes
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 88 // DW_AT_call_file
.b8 11 // DW_FORM_data1
.b8 89 // DW_AT_call_line
.b8 11 // DW_FORM_data1
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 6 // Abbreviation Code
.b8 29 // DW_TAG_inlined_subroutine
.b8 0 // DW_CHILDREN_no
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 88 // DW_AT_call_file
.b8 11 // DW_FORM_data1
.b8 89 // DW_AT_call_line
.b8 11 // DW_FORM_data1
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 7 // Abbreviation Code
.b8 29 // DW_TAG_inlined_subroutine
.b8 0 // DW_CHILDREN_no
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 88 // DW_AT_call_file
.b8 11 // DW_FORM_data1
.b8 89 // DW_AT_call_line
.b8 5 // DW_FORM_data2
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 8 // Abbreviation Code
.b8 29 // DW_TAG_inlined_subroutine
.b8 1 // DW_CHILDREN_yes
.b8 49 // DW_AT_abstract_origin
.b8 19 // DW_FORM_ref4
.b8 17 // DW_AT_low_pc
.b8 1 // DW_FORM_addr
.b8 18 // DW_AT_high_pc
.b8 1 // DW_FORM_addr
.b8 88 // DW_AT_call_file
.b8 11 // DW_FORM_data1
.b8 89 // DW_AT_call_line
.b8 5 // DW_FORM_data2
.b8 0 // EOM(1)
.b8 0 // EOM(2)
.b8 0 // EOM(3)
}
.section .debug_info
{
.b32 44 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 1 // Abbrev [1] 0xb:0x25 DW_TAG_compile_unit
.b8 106 // DW_AT_producer
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b8 31 // DW_AT_language
.b8 0
.b8 106 // DW_AT_name
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 46 // DW_AT_comp_dir
.b8 0
.b64 $L__func_begin0 // DW_AT_low_pc
.b64 $L__func_end0 // DW_AT_high_pc
.b32 3029 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 2 // Abbrev [2] 0xb:0xbce DW_TAG_compile_unit
.b8 106 // DW_AT_producer
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b8 31 // DW_AT_language
.b8 0
.b8 106 // DW_AT_name
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 46 // DW_AT_comp_dir
.b8 0
.b64 $L__func_begin1 // DW_AT_low_pc
.b64 $L__func_end1 // DW_AT_high_pc
.b8 3 // Abbrev [3] 0x30:0x5 DW_TAG_subprogram
.b8 43 // DW_AT_name
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x35:0xf DW_TAG_subprogram
.b8 103 // DW_AT_name
.b8 101
.b8 116
.b8 112
.b8 114
.b8 111
.b8 112
.b8 101
.b8 114
.b8 116
.b8 121
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x44:0x10 DW_TAG_subprogram
.b8 108 // DW_AT_name
.b8 111
.b8 99
.b8 107
.b8 95
.b8 111
.b8 117
.b8 116
.b8 112
.b8 117
.b8 116
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x54:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x67:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x7a:0x13 DW_TAG_subprogram
.b8 108 // DW_AT_name
.b8 108
.b8 118
.b8 109
.b8 95
.b8 97
.b8 116
.b8 111
.b8 109
.b8 105
.b8 99
.b8 95
.b8 99
.b8 97
.b8 115
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x8d:0xf DW_TAG_subprogram
.b8 97 // DW_AT_name
.b8 116
.b8 111
.b8 109
.b8 105
.b8 99
.b8 95
.b8 99
.b8 97
.b8 115
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x9c:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xaf:0xa DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 105
.b8 110
.b8 100
.b8 101
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xb9:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 95
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xc8:0xe DW_TAG_subprogram
.b8 35 // DW_AT_name
.b8 116
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xd6:0x5 DW_TAG_subprogram
.b8 43 // DW_AT_name
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xdb:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 95
.b8 121
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xea:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 95
.b8 122
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xf9:0xe DW_TAG_subprogram
.b8 98 // DW_AT_name
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 95
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x107:0xd DW_TAG_subprogram
.b8 35 // DW_AT_name
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x114:0xe DW_TAG_subprogram
.b8 98 // DW_AT_name
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 95
.b8 121
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x122:0xe DW_TAG_subprogram
.b8 98 // DW_AT_name
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 95
.b8 122
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x130:0x11 DW_TAG_subprogram
.b8 117 // DW_AT_name
.b8 110
.b8 115
.b8 97
.b8 102
.b8 101
.b8 95
.b8 115
.b8 116
.b8 111
.b8 114
.b8 101
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x141:0x10 DW_TAG_subprogram
.b8 115 // DW_AT_name
.b8 101
.b8 116
.b8 112
.b8 114
.b8 111
.b8 112
.b8 101
.b8 114
.b8 116
.b8 121
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x151:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 102
.b8 101
.b8 110
.b8 99
.b8 101
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x160:0xf DW_TAG_subprogram
.b8 117 // DW_AT_name
.b8 110
.b8 115
.b8 97
.b8 102
.b8 101
.b8 95
.b8 108
.b8 111
.b8 97
.b8 100
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x16f:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x175:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x17b:0x6 DW_TAG_subprogram
.b8 33 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x181:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x194:0xd DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 99
.b8 117
.b8 112
.b8 114
.b8 105
.b8 110
.b8 116
.b8 102
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1a1:0x7 DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 101
.b8 113
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1a8:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1ae:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 4 // Abbrev [4] 0x1b4:0xa24 DW_TAG_subprogram
.b64 $L__func_begin1 // DW_AT_low_pc
.b64 $L__func_end1 // DW_AT_high_pc
.b8 114 // DW_AT_name
.b8 101
.b8 112
.b8 111
.b8 114
.b8 116
.b8 95
.b8 101
.b8 120
.b8 99
.b8 101
.b8 112
.b8 116
.b8 105
.b8 111
.b8 110
.b8 0
.b8 5 // Abbrev [5] 0x1d6:0x80c DW_TAG_inlined_subroutine
.b32 68 // DW_AT_abstract_origin
.b64 $L__tmp2 // DW_AT_low_pc
.b64 $L__tmp55 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 146 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x1ed:0x2f DW_TAG_inlined_subroutine
.b32 53 // DW_AT_abstract_origin
.b64 $L__tmp2 // DW_AT_low_pc
.b64 $L__tmp3 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 129 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x204:0x17 DW_TAG_inlined_subroutine
.b32 48 // DW_AT_abstract_origin
.b64 $L__tmp2 // DW_AT_low_pc
.b64 $L__tmp3 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 64 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x21c:0x5f DW_TAG_inlined_subroutine
.b32 141 // DW_AT_abstract_origin
.b64 $L__tmp3 // DW_AT_low_pc
.b64 $L__tmp4 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 129 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x233:0x47 DW_TAG_inlined_subroutine
.b32 122 // DW_AT_abstract_origin
.b64 $L__tmp3 // DW_AT_low_pc
.b64 $L__tmp4 // DW_AT_high_pc
.b8 9 // DW_AT_call_file
.b8 139 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x24a:0x2f DW_TAG_inlined_subroutine
.b32 103 // DW_AT_abstract_origin
.b64 $L__tmp3 // DW_AT_low_pc
.b64 $L__tmp4 // DW_AT_high_pc
.b8 9 // DW_AT_call_file
.b8 110 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x261:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp3 // DW_AT_low_pc
.b64 $L__tmp4 // DW_AT_high_pc
.b8 9 // DW_AT_call_file
.b8 110 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x27b:0x17a DW_TAG_inlined_subroutine
.b32 200 // DW_AT_abstract_origin
.b64 $L__tmp5 // DW_AT_low_pc
.b64 $L__tmp11 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 131 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x292:0x76 DW_TAG_inlined_subroutine
.b32 185 // DW_AT_abstract_origin
.b64 $L__tmp5 // DW_AT_low_pc
.b64 $L__tmp7 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x2a9:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp5 // DW_AT_low_pc
.b64 $L__tmp6 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x2c0:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp5 // DW_AT_low_pc
.b64 $L__tmp6 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x2d7:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp5 // DW_AT_low_pc
.b64 $L__tmp6 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x2f0:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp6 // DW_AT_low_pc
.b64 $L__tmp7 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x308:0x76 DW_TAG_inlined_subroutine
.b32 219 // DW_AT_abstract_origin
.b64 $L__tmp7 // DW_AT_low_pc
.b64 $L__tmp9 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x31f:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp7 // DW_AT_low_pc
.b64 $L__tmp8 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x336:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp7 // DW_AT_low_pc
.b64 $L__tmp8 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x34d:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp7 // DW_AT_low_pc
.b64 $L__tmp8 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x366:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp8 // DW_AT_low_pc
.b64 $L__tmp9 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x37e:0x76 DW_TAG_inlined_subroutine
.b32 234 // DW_AT_abstract_origin
.b64 $L__tmp9 // DW_AT_low_pc
.b64 $L__tmp11 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x395:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp9 // DW_AT_low_pc
.b64 $L__tmp10 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x3ac:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp9 // DW_AT_low_pc
.b64 $L__tmp10 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x3c3:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp9 // DW_AT_low_pc
.b64 $L__tmp10 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x3dc:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp10 // DW_AT_low_pc
.b64 $L__tmp11 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x3f5:0x17a DW_TAG_inlined_subroutine
.b32 263 // DW_AT_abstract_origin
.b64 $L__tmp11 // DW_AT_low_pc
.b64 $L__tmp17 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 131 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x40c:0x76 DW_TAG_inlined_subroutine
.b32 249 // DW_AT_abstract_origin
.b64 $L__tmp11 // DW_AT_low_pc
.b64 $L__tmp13 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x423:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp11 // DW_AT_low_pc
.b64 $L__tmp12 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x43a:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp11 // DW_AT_low_pc
.b64 $L__tmp12 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x451:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp11 // DW_AT_low_pc
.b64 $L__tmp12 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x46a:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp12 // DW_AT_low_pc
.b64 $L__tmp13 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x482:0x76 DW_TAG_inlined_subroutine
.b32 276 // DW_AT_abstract_origin
.b64 $L__tmp13 // DW_AT_low_pc
.b64 $L__tmp15 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x499:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp13 // DW_AT_low_pc
.b64 $L__tmp14 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x4b0:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp13 // DW_AT_low_pc
.b64 $L__tmp14 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x4c7:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp13 // DW_AT_low_pc
.b64 $L__tmp14 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x4e0:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp14 // DW_AT_low_pc
.b64 $L__tmp15 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x4f8:0x76 DW_TAG_inlined_subroutine
.b32 290 // DW_AT_abstract_origin
.b64 $L__tmp15 // DW_AT_low_pc
.b64 $L__tmp17 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x50f:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp15 // DW_AT_low_pc
.b64 $L__tmp16 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x526:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp15 // DW_AT_low_pc
.b64 $L__tmp16 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x53d:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp15 // DW_AT_low_pc
.b64 $L__tmp16 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x556:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp16 // DW_AT_low_pc
.b64 $L__tmp17 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x56f:0x76 DW_TAG_inlined_subroutine
.b32 321 // DW_AT_abstract_origin
.b64 $L__tmp17 // DW_AT_low_pc
.b64 $L__tmp19 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 131 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x586:0x2f DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp17 // DW_AT_low_pc
.b64 $L__tmp18 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 83 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x59d:0x17 DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp17 // DW_AT_low_pc
.b64 $L__tmp18 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 180 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x5b5:0x2f DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp18 // DW_AT_low_pc
.b64 $L__tmp19 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 85 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x5cc:0x17 DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp18 // DW_AT_low_pc
.b64 $L__tmp19 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 180 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x5e5:0x17 DW_TAG_inlined_subroutine
.b32 337 // DW_AT_abstract_origin
.b64 $L__tmp19 // DW_AT_low_pc
.b64 $L__tmp20 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 132 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x5fc:0xa5 DW_TAG_inlined_subroutine
.b32 53 // DW_AT_abstract_origin
.b64 $L__tmp25 // DW_AT_low_pc
.b64 $L__tmp52 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x613:0x2f DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp25 // DW_AT_low_pc
.b64 $L__tmp26 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 62 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x62a:0x17 DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp25 // DW_AT_low_pc
.b64 $L__tmp26 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x642:0x2f DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp28 // DW_AT_low_pc
.b64 $L__tmp38 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 66 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x659:0x17 DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp28 // DW_AT_low_pc
.b64 $L__tmp38 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x671:0x2f DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp41 // DW_AT_low_pc
.b64 $L__tmp52 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 68 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x688:0x17 DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp41 // DW_AT_low_pc
.b64 $L__tmp52 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x6a1:0x30 DW_TAG_inlined_subroutine
.b32 367 // DW_AT_abstract_origin
.b64 $L__tmp26 // DW_AT_low_pc
.b64 $L__tmp27 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 7 // Abbrev [7] 0x6b8:0x18 DW_TAG_inlined_subroutine
.b32 367 // DW_AT_abstract_origin
.b64 $L__tmp26 // DW_AT_low_pc
.b64 $L__tmp27 // DW_AT_high_pc
.b8 7 // DW_AT_call_file
.b8 227 // DW_AT_call_line
.b8 1
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x6d1:0x104 DW_TAG_inlined_subroutine
.b32 200 // DW_AT_abstract_origin
.b64 $L__tmp29 // DW_AT_low_pc
.b64 $L__tmp37 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x6e8:0x76 DW_TAG_inlined_subroutine
.b32 185 // DW_AT_abstract_origin
.b64 $L__tmp29 // DW_AT_low_pc
.b64 $L__tmp31 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x6ff:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp29 // DW_AT_low_pc
.b64 $L__tmp30 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x716:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp29 // DW_AT_low_pc
.b64 $L__tmp30 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x72d:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp29 // DW_AT_low_pc
.b64 $L__tmp30 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x746:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp30 // DW_AT_low_pc
.b64 $L__tmp31 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x75e:0x76 DW_TAG_inlined_subroutine
.b32 219 // DW_AT_abstract_origin
.b64 $L__tmp35 // DW_AT_low_pc
.b64 $L__tmp37 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x775:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp35 // DW_AT_low_pc
.b64 $L__tmp36 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x78c:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp35 // DW_AT_low_pc
.b64 $L__tmp36 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x7a3:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp35 // DW_AT_low_pc
.b64 $L__tmp36 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x7bc:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp36 // DW_AT_low_pc
.b64 $L__tmp37 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x7d5:0x92 DW_TAG_inlined_subroutine
.b32 430 // DW_AT_abstract_origin
.b64 $L__tmp31 // DW_AT_low_pc
.b64 $L__tmp54 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x7ec:0x7a DW_TAG_inlined_subroutine
.b32 424 // DW_AT_abstract_origin
.b64 $L__tmp31 // DW_AT_low_pc
.b64 $L__tmp54 // DW_AT_high_pc
.b8 11 // DW_AT_call_file
.b8 244 // DW_AT_call_line
.b8 8 // Abbrev [8] 0x803:0x62 DW_TAG_inlined_subroutine
.b32 417 // DW_AT_abstract_origin
.b64 $L__tmp31 // DW_AT_low_pc
.b64 $L__tmp54 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 35 // DW_AT_call_line
.b8 2
.b8 7 // Abbrev [7] 0x81b:0x18 DW_TAG_inlined_subroutine
.b32 367 // DW_AT_abstract_origin
.b64 $L__tmp31 // DW_AT_low_pc
.b64 $L__tmp45 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 39 // DW_AT_call_line
.b8 2
.b8 8 // Abbrev [8] 0x833:0x31 DW_TAG_inlined_subroutine
.b32 417 // DW_AT_abstract_origin
.b64 $L__tmp38 // DW_AT_low_pc
.b64 $L__tmp54 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 45 // DW_AT_call_line
.b8 2
.b8 7 // Abbrev [7] 0x84b:0x18 DW_TAG_inlined_subroutine
.b32 367 // DW_AT_abstract_origin
.b64 $L__tmp38 // DW_AT_low_pc
.b64 $L__tmp53 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 39 // DW_AT_call_line
.b8 2
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x867:0x17a DW_TAG_inlined_subroutine
.b32 263 // DW_AT_abstract_origin
.b64 $L__tmp42 // DW_AT_low_pc
.b64 $L__tmp51 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x87e:0x76 DW_TAG_inlined_subroutine
.b32 249 // DW_AT_abstract_origin
.b64 $L__tmp42 // DW_AT_low_pc
.b64 $L__tmp44 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x895:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp42 // DW_AT_low_pc
.b64 $L__tmp43 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x8ac:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp42 // DW_AT_low_pc
.b64 $L__tmp43 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x8c3:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp42 // DW_AT_low_pc
.b64 $L__tmp43 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x8dc:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp43 // DW_AT_low_pc
.b64 $L__tmp44 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x8f4:0x76 DW_TAG_inlined_subroutine
.b32 290 // DW_AT_abstract_origin
.b64 $L__tmp46 // DW_AT_low_pc
.b64 $L__tmp48 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x90b:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp46 // DW_AT_low_pc
.b64 $L__tmp47 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x922:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp46 // DW_AT_low_pc
.b64 $L__tmp47 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x939:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp46 // DW_AT_low_pc
.b64 $L__tmp47 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x952:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp47 // DW_AT_low_pc
.b64 $L__tmp48 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x96a:0x76 DW_TAG_inlined_subroutine
.b32 276 // DW_AT_abstract_origin
.b64 $L__tmp49 // DW_AT_low_pc
.b64 $L__tmp51 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x981:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp49 // DW_AT_low_pc
.b64 $L__tmp50 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x998:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp49 // DW_AT_low_pc
.b64 $L__tmp50 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x9af:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp49 // DW_AT_low_pc
.b64 $L__tmp50 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x9c8:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp50 // DW_AT_low_pc
.b64 $L__tmp51 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x9e2:0x47 DW_TAG_inlined_subroutine
.b32 53 // DW_AT_abstract_origin
.b64 $L__tmp21 // DW_AT_low_pc
.b64 $L__tmp22 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 148 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x9f9:0x2f DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp21 // DW_AT_low_pc
.b64 $L__tmp22 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 70 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xa10:0x17 DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp21 // DW_AT_low_pc
.b64 $L__tmp22 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xa29:0x49 DW_TAG_inlined_subroutine
.b32 379 // DW_AT_abstract_origin
.b64 $L__tmp22 // DW_AT_low_pc
.b64 $L__tmp23 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 148 // DW_AT_call_line
.b8 8 // Abbrev [8] 0xa40:0x31 DW_TAG_inlined_subroutine
.b32 373 // DW_AT_abstract_origin
.b64 $L__tmp22 // DW_AT_low_pc
.b64 $L__tmp23 // DW_AT_high_pc
.b8 12 // DW_AT_call_file
.b8 21 // DW_AT_call_line
.b8 1
.b8 7 // Abbrev [7] 0xa58:0x18 DW_TAG_inlined_subroutine
.b32 367 // DW_AT_abstract_origin
.b64 $L__tmp22 // DW_AT_low_pc
.b64 $L__tmp23 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 1
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xa72:0x47 DW_TAG_inlined_subroutine
.b32 404 // DW_AT_abstract_origin
.b64 $L__tmp24 // DW_AT_low_pc
.b64 $L__tmp56 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 151 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xa89:0x2f DW_TAG_inlined_subroutine
.b32 385 // DW_AT_abstract_origin
.b64 $L__tmp24 // DW_AT_low_pc
.b64 $L__tmp56 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xaa0:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp24 // DW_AT_low_pc
.b64 $L__tmp56 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xab9:0x47 DW_TAG_inlined_subroutine
.b32 53 // DW_AT_abstract_origin
.b64 $L__tmp56 // DW_AT_low_pc
.b64 $L__tmp57 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xad0:0x2f DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp56 // DW_AT_low_pc
.b64 $L__tmp57 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 72 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xae7:0x17 DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp56 // DW_AT_low_pc
.b64 $L__tmp57 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xb00:0x49 DW_TAG_inlined_subroutine
.b32 379 // DW_AT_abstract_origin
.b64 $L__tmp57 // DW_AT_low_pc
.b64 $L__tmp58 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 8 // Abbrev [8] 0xb17:0x31 DW_TAG_inlined_subroutine
.b32 373 // DW_AT_abstract_origin
.b64 $L__tmp57 // DW_AT_low_pc
.b64 $L__tmp58 // DW_AT_high_pc
.b8 12 // DW_AT_call_file
.b8 21 // DW_AT_call_line
.b8 1
.b8 7 // Abbrev [7] 0xb2f:0x18 DW_TAG_inlined_subroutine
.b32 367 // DW_AT_abstract_origin
.b64 $L__tmp57 // DW_AT_low_pc
.b64 $L__tmp58 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 1
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xb49:0x47 DW_TAG_inlined_subroutine
.b32 404 // DW_AT_abstract_origin
.b64 $L__tmp59 // DW_AT_low_pc
.b64 $L__tmp60 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 154 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xb60:0x2f DW_TAG_inlined_subroutine
.b32 385 // DW_AT_abstract_origin
.b64 $L__tmp59 // DW_AT_low_pc
.b64 $L__tmp60 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xb77:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp59 // DW_AT_low_pc
.b64 $L__tmp60 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xb90:0x47 DW_TAG_inlined_subroutine
.b32 404 // DW_AT_abstract_origin
.b64 $L__tmp60 // DW_AT_low_pc
.b64 $L__tmp61 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 156 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xba7:0x2f DW_TAG_inlined_subroutine
.b32 385 // DW_AT_abstract_origin
.b64 $L__tmp60 // DW_AT_low_pc
.b64 $L__tmp61 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xbbe:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp60 // DW_AT_low_pc
.b64 $L__tmp61 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b32 100 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 2 // Abbrev [2] 0xb:0x5d DW_TAG_compile_unit
.b8 106 // DW_AT_producer
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b8 31 // DW_AT_language
.b8 0
.b8 106 // DW_AT_name
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 46 // DW_AT_comp_dir
.b8 0
.b64 $L__func_begin2 // DW_AT_low_pc
.b64 $L__func_end2 // DW_AT_high_pc
.b8 3 // Abbrev [3] 0x30:0xc DW_TAG_subprogram
.b8 103 // DW_AT_name
.b8 101
.b8 116
.b8 105
.b8 110
.b8 100
.b8 101
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 4 // Abbrev [4] 0x3c:0x2b DW_TAG_subprogram
.b64 $L__func_begin2 // DW_AT_low_pc
.b64 $L__func_end2 // DW_AT_high_pc
.b8 102 // DW_AT_name
.b8 0
.b8 6 // Abbrev [6] 0x4f:0x17 DW_TAG_inlined_subroutine
.b32 48 // DW_AT_abstract_origin
.b64 $L__tmp63 // DW_AT_low_pc
.b64 $L__tmp64 // DW_AT_high_pc
.b8 14 // DW_AT_call_file
.b8 5 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b32 44 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 1 // Abbrev [1] 0xb:0x25 DW_TAG_compile_unit
.b8 106 // DW_AT_producer
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b8 31 // DW_AT_language
.b8 0
.b8 106 // DW_AT_name
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 46 // DW_AT_comp_dir
.b8 0
.b64 $L__func_begin3 // DW_AT_low_pc
.b64 $L__func_end3 // DW_AT_high_pc
.b32 3026 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 2 // Abbrev [2] 0xb:0xbcb DW_TAG_compile_unit
.b8 106 // DW_AT_producer
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b8 31 // DW_AT_language
.b8 0
.b8 106 // DW_AT_name
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 46 // DW_AT_comp_dir
.b8 0
.b64 $L__func_begin4 // DW_AT_low_pc
.b64 $L__func_end4 // DW_AT_high_pc
.b8 3 // Abbrev [3] 0x30:0x5 DW_TAG_subprogram
.b8 43 // DW_AT_name
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x35:0xf DW_TAG_subprogram
.b8 103 // DW_AT_name
.b8 101
.b8 116
.b8 112
.b8 114
.b8 111
.b8 112
.b8 101
.b8 114
.b8 116
.b8 121
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x44:0x10 DW_TAG_subprogram
.b8 108 // DW_AT_name
.b8 111
.b8 99
.b8 107
.b8 95
.b8 111
.b8 117
.b8 116
.b8 112
.b8 117
.b8 116
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x54:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x67:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x7a:0x13 DW_TAG_subprogram
.b8 108 // DW_AT_name
.b8 108
.b8 118
.b8 109
.b8 95
.b8 97
.b8 116
.b8 111
.b8 109
.b8 105
.b8 99
.b8 95
.b8 99
.b8 97
.b8 115
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x8d:0xf DW_TAG_subprogram
.b8 97 // DW_AT_name
.b8 116
.b8 111
.b8 109
.b8 105
.b8 99
.b8 95
.b8 99
.b8 97
.b8 115
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x9c:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xaf:0xa DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 105
.b8 110
.b8 100
.b8 101
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xb9:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 95
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xc8:0xe DW_TAG_subprogram
.b8 35 // DW_AT_name
.b8 116
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xd6:0x5 DW_TAG_subprogram
.b8 43 // DW_AT_name
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xdb:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 95
.b8 121
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xea:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 73
.b8 100
.b8 120
.b8 95
.b8 122
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0xf9:0xe DW_TAG_subprogram
.b8 98 // DW_AT_name
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 95
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x107:0xd DW_TAG_subprogram
.b8 35 // DW_AT_name
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x114:0xe DW_TAG_subprogram
.b8 98 // DW_AT_name
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 95
.b8 121
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x122:0xe DW_TAG_subprogram
.b8 98 // DW_AT_name
.b8 108
.b8 111
.b8 99
.b8 107
.b8 73
.b8 100
.b8 120
.b8 95
.b8 122
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x130:0x11 DW_TAG_subprogram
.b8 117 // DW_AT_name
.b8 110
.b8 115
.b8 97
.b8 102
.b8 101
.b8 95
.b8 115
.b8 116
.b8 111
.b8 114
.b8 101
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x141:0x10 DW_TAG_subprogram
.b8 115 // DW_AT_name
.b8 101
.b8 116
.b8 112
.b8 114
.b8 111
.b8 112
.b8 101
.b8 114
.b8 116
.b8 121
.b8 33
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x151:0xf DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 102
.b8 101
.b8 110
.b8 99
.b8 101
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x160:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x173:0xd DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 99
.b8 117
.b8 112
.b8 114
.b8 105
.b8 110
.b8 116
.b8 102
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x180:0xf DW_TAG_subprogram
.b8 117 // DW_AT_name
.b8 110
.b8 115
.b8 97
.b8 102
.b8 101
.b8 95
.b8 108
.b8 111
.b8 97
.b8 100
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x18f:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x195:0x7 DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 101
.b8 113
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x19c:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1a2:0x6 DW_TAG_subprogram
.b8 61 // DW_AT_name
.b8 61
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1a8:0x16 DW_TAG_subprogram
.b8 116 // DW_AT_name
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 102
.b8 101
.b8 110
.b8 99
.b8 101
.b8 95
.b8 115
.b8 121
.b8 115
.b8 116
.b8 101
.b8 109
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1be:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1d1:0xc DW_TAG_subprogram
.b8 95 // DW_AT_name
.b8 97
.b8 115
.b8 109
.b8 99
.b8 97
.b8 108
.b8 108
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x1dd:0x8 DW_TAG_subprogram
.b8 101 // DW_AT_name
.b8 120
.b8 105
.b8 116
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 4 // Abbrev [4] 0x1e5:0x9f0 DW_TAG_subprogram
.b64 $L__func_begin4 // DW_AT_low_pc
.b64 $L__func_end4 // DW_AT_high_pc
.b8 115 // DW_AT_name
.b8 105
.b8 103
.b8 110
.b8 97
.b8 108
.b8 95
.b8 101
.b8 120
.b8 99
.b8 101
.b8 112
.b8 116
.b8 105
.b8 111
.b8 110
.b8 0
.b8 5 // Abbrev [5] 0x207:0x882 DW_TAG_inlined_subroutine
.b32 68 // DW_AT_abstract_origin
.b64 $L__tmp67 // DW_AT_low_pc
.b64 $L__tmp118 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 193 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x21e:0x2f DW_TAG_inlined_subroutine
.b32 53 // DW_AT_abstract_origin
.b64 $L__tmp67 // DW_AT_low_pc
.b64 $L__tmp68 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 129 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x235:0x17 DW_TAG_inlined_subroutine
.b32 48 // DW_AT_abstract_origin
.b64 $L__tmp67 // DW_AT_low_pc
.b64 $L__tmp68 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 64 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x24d:0x5f DW_TAG_inlined_subroutine
.b32 141 // DW_AT_abstract_origin
.b64 $L__tmp68 // DW_AT_low_pc
.b64 $L__tmp69 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 129 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x264:0x47 DW_TAG_inlined_subroutine
.b32 122 // DW_AT_abstract_origin
.b64 $L__tmp68 // DW_AT_low_pc
.b64 $L__tmp69 // DW_AT_high_pc
.b8 9 // DW_AT_call_file
.b8 139 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x27b:0x2f DW_TAG_inlined_subroutine
.b32 103 // DW_AT_abstract_origin
.b64 $L__tmp68 // DW_AT_low_pc
.b64 $L__tmp69 // DW_AT_high_pc
.b8 9 // DW_AT_call_file
.b8 110 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x292:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp68 // DW_AT_low_pc
.b64 $L__tmp69 // DW_AT_high_pc
.b8 9 // DW_AT_call_file
.b8 110 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x2ac:0x17a DW_TAG_inlined_subroutine
.b32 200 // DW_AT_abstract_origin
.b64 $L__tmp70 // DW_AT_low_pc
.b64 $L__tmp76 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 131 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x2c3:0x76 DW_TAG_inlined_subroutine
.b32 185 // DW_AT_abstract_origin
.b64 $L__tmp70 // DW_AT_low_pc
.b64 $L__tmp72 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x2da:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp70 // DW_AT_low_pc
.b64 $L__tmp71 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x2f1:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp70 // DW_AT_low_pc
.b64 $L__tmp71 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x308:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp70 // DW_AT_low_pc
.b64 $L__tmp71 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x321:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp71 // DW_AT_low_pc
.b64 $L__tmp72 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x339:0x76 DW_TAG_inlined_subroutine
.b32 219 // DW_AT_abstract_origin
.b64 $L__tmp72 // DW_AT_low_pc
.b64 $L__tmp74 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x350:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp72 // DW_AT_low_pc
.b64 $L__tmp73 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x367:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp72 // DW_AT_low_pc
.b64 $L__tmp73 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x37e:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp72 // DW_AT_low_pc
.b64 $L__tmp73 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x397:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp73 // DW_AT_low_pc
.b64 $L__tmp74 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x3af:0x76 DW_TAG_inlined_subroutine
.b32 234 // DW_AT_abstract_origin
.b64 $L__tmp74 // DW_AT_low_pc
.b64 $L__tmp76 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x3c6:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp74 // DW_AT_low_pc
.b64 $L__tmp75 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x3dd:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp74 // DW_AT_low_pc
.b64 $L__tmp75 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x3f4:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp74 // DW_AT_low_pc
.b64 $L__tmp75 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x40d:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp75 // DW_AT_low_pc
.b64 $L__tmp76 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x426:0x17a DW_TAG_inlined_subroutine
.b32 263 // DW_AT_abstract_origin
.b64 $L__tmp76 // DW_AT_low_pc
.b64 $L__tmp82 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 131 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x43d:0x76 DW_TAG_inlined_subroutine
.b32 249 // DW_AT_abstract_origin
.b64 $L__tmp76 // DW_AT_low_pc
.b64 $L__tmp78 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x454:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp76 // DW_AT_low_pc
.b64 $L__tmp77 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x46b:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp76 // DW_AT_low_pc
.b64 $L__tmp77 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x482:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp76 // DW_AT_low_pc
.b64 $L__tmp77 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x49b:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp77 // DW_AT_low_pc
.b64 $L__tmp78 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x4b3:0x76 DW_TAG_inlined_subroutine
.b32 276 // DW_AT_abstract_origin
.b64 $L__tmp78 // DW_AT_low_pc
.b64 $L__tmp80 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x4ca:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp78 // DW_AT_low_pc
.b64 $L__tmp79 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x4e1:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp78 // DW_AT_low_pc
.b64 $L__tmp79 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x4f8:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp78 // DW_AT_low_pc
.b64 $L__tmp79 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x511:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp79 // DW_AT_low_pc
.b64 $L__tmp80 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x529:0x76 DW_TAG_inlined_subroutine
.b32 290 // DW_AT_abstract_origin
.b64 $L__tmp80 // DW_AT_low_pc
.b64 $L__tmp82 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x540:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp80 // DW_AT_low_pc
.b64 $L__tmp81 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x557:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp80 // DW_AT_low_pc
.b64 $L__tmp81 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x56e:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp80 // DW_AT_low_pc
.b64 $L__tmp81 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x587:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp81 // DW_AT_low_pc
.b64 $L__tmp82 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x5a0:0x76 DW_TAG_inlined_subroutine
.b32 321 // DW_AT_abstract_origin
.b64 $L__tmp82 // DW_AT_low_pc
.b64 $L__tmp84 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 131 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x5b7:0x2f DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp82 // DW_AT_low_pc
.b64 $L__tmp83 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 83 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x5ce:0x17 DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp82 // DW_AT_low_pc
.b64 $L__tmp83 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 180 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x5e6:0x2f DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp83 // DW_AT_low_pc
.b64 $L__tmp84 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 85 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x5fd:0x17 DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp83 // DW_AT_low_pc
.b64 $L__tmp84 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 180 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x616:0x17 DW_TAG_inlined_subroutine
.b32 337 // DW_AT_abstract_origin
.b64 $L__tmp84 // DW_AT_low_pc
.b64 $L__tmp85 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 132 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x62d:0xa5 DW_TAG_inlined_subroutine
.b32 53 // DW_AT_abstract_origin
.b64 $L__tmp87 // DW_AT_low_pc
.b64 $L__tmp115 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x644:0x2f DW_TAG_inlined_subroutine
.b32 384 // DW_AT_abstract_origin
.b64 $L__tmp87 // DW_AT_low_pc
.b64 $L__tmp88 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 62 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x65b:0x17 DW_TAG_inlined_subroutine
.b32 384 // DW_AT_abstract_origin
.b64 $L__tmp87 // DW_AT_low_pc
.b64 $L__tmp88 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x673:0x2f DW_TAG_inlined_subroutine
.b32 384 // DW_AT_abstract_origin
.b64 $L__tmp90 // DW_AT_low_pc
.b64 $L__tmp101 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 66 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x68a:0x17 DW_TAG_inlined_subroutine
.b32 384 // DW_AT_abstract_origin
.b64 $L__tmp90 // DW_AT_low_pc
.b64 $L__tmp101 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x6a2:0x2f DW_TAG_inlined_subroutine
.b32 384 // DW_AT_abstract_origin
.b64 $L__tmp104 // DW_AT_low_pc
.b64 $L__tmp115 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 68 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x6b9:0x17 DW_TAG_inlined_subroutine
.b32 384 // DW_AT_abstract_origin
.b64 $L__tmp104 // DW_AT_low_pc
.b64 $L__tmp115 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 153 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x6d2:0x30 DW_TAG_inlined_subroutine
.b32 399 // DW_AT_abstract_origin
.b64 $L__tmp88 // DW_AT_low_pc
.b64 $L__tmp89 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 7 // Abbrev [7] 0x6e9:0x18 DW_TAG_inlined_subroutine
.b32 399 // DW_AT_abstract_origin
.b64 $L__tmp88 // DW_AT_low_pc
.b64 $L__tmp89 // DW_AT_high_pc
.b8 7 // DW_AT_call_file
.b8 227 // DW_AT_call_line
.b8 1
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x702:0x17a DW_TAG_inlined_subroutine
.b32 200 // DW_AT_abstract_origin
.b64 $L__tmp91 // DW_AT_low_pc
.b64 $L__tmp100 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x719:0x76 DW_TAG_inlined_subroutine
.b32 185 // DW_AT_abstract_origin
.b64 $L__tmp91 // DW_AT_low_pc
.b64 $L__tmp93 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x730:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp91 // DW_AT_low_pc
.b64 $L__tmp92 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x747:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp91 // DW_AT_low_pc
.b64 $L__tmp92 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x75e:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp91 // DW_AT_low_pc
.b64 $L__tmp92 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x777:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp92 // DW_AT_low_pc
.b64 $L__tmp93 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x78f:0x76 DW_TAG_inlined_subroutine
.b32 234 // DW_AT_abstract_origin
.b64 $L__tmp95 // DW_AT_low_pc
.b64 $L__tmp97 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x7a6:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp95 // DW_AT_low_pc
.b64 $L__tmp96 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x7bd:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp95 // DW_AT_low_pc
.b64 $L__tmp96 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x7d4:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp95 // DW_AT_low_pc
.b64 $L__tmp96 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x7ed:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp96 // DW_AT_low_pc
.b64 $L__tmp97 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x805:0x76 DW_TAG_inlined_subroutine
.b32 219 // DW_AT_abstract_origin
.b64 $L__tmp98 // DW_AT_low_pc
.b64 $L__tmp100 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 92 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x81c:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp98 // DW_AT_low_pc
.b64 $L__tmp99 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x833:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp98 // DW_AT_low_pc
.b64 $L__tmp99 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x84a:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp98 // DW_AT_low_pc
.b64 $L__tmp99 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x863:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp99 // DW_AT_low_pc
.b64 $L__tmp100 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 46 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x87c:0x92 DW_TAG_inlined_subroutine
.b32 418 // DW_AT_abstract_origin
.b64 $L__tmp93 // DW_AT_low_pc
.b64 $L__tmp117 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x893:0x7a DW_TAG_inlined_subroutine
.b32 412 // DW_AT_abstract_origin
.b64 $L__tmp93 // DW_AT_low_pc
.b64 $L__tmp117 // DW_AT_high_pc
.b8 11 // DW_AT_call_file
.b8 244 // DW_AT_call_line
.b8 8 // Abbrev [8] 0x8aa:0x62 DW_TAG_inlined_subroutine
.b32 405 // DW_AT_abstract_origin
.b64 $L__tmp93 // DW_AT_low_pc
.b64 $L__tmp117 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 35 // DW_AT_call_line
.b8 2
.b8 7 // Abbrev [7] 0x8c2:0x18 DW_TAG_inlined_subroutine
.b32 399 // DW_AT_abstract_origin
.b64 $L__tmp93 // DW_AT_low_pc
.b64 $L__tmp108 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 39 // DW_AT_call_line
.b8 2
.b8 8 // Abbrev [8] 0x8da:0x31 DW_TAG_inlined_subroutine
.b32 405 // DW_AT_abstract_origin
.b64 $L__tmp101 // DW_AT_low_pc
.b64 $L__tmp117 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 45 // DW_AT_call_line
.b8 2
.b8 7 // Abbrev [7] 0x8f2:0x18 DW_TAG_inlined_subroutine
.b32 399 // DW_AT_abstract_origin
.b64 $L__tmp101 // DW_AT_low_pc
.b64 $L__tmp116 // DW_AT_high_pc
.b8 8 // DW_AT_call_file
.b8 39 // DW_AT_call_line
.b8 2
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x90e:0x17a DW_TAG_inlined_subroutine
.b32 263 // DW_AT_abstract_origin
.b64 $L__tmp105 // DW_AT_low_pc
.b64 $L__tmp114 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 134 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x925:0x76 DW_TAG_inlined_subroutine
.b32 249 // DW_AT_abstract_origin
.b64 $L__tmp105 // DW_AT_low_pc
.b64 $L__tmp107 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x93c:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp105 // DW_AT_low_pc
.b64 $L__tmp106 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x953:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp105 // DW_AT_low_pc
.b64 $L__tmp106 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x96a:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp105 // DW_AT_low_pc
.b64 $L__tmp106 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x983:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp106 // DW_AT_low_pc
.b64 $L__tmp107 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0x99b:0x76 DW_TAG_inlined_subroutine
.b32 290 // DW_AT_abstract_origin
.b64 $L__tmp109 // DW_AT_low_pc
.b64 $L__tmp111 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x9b2:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp109 // DW_AT_low_pc
.b64 $L__tmp110 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0x9c9:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp109 // DW_AT_low_pc
.b64 $L__tmp110 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0x9e0:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp109 // DW_AT_low_pc
.b64 $L__tmp110 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0x9f9:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp110 // DW_AT_low_pc
.b64 $L__tmp111 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xa11:0x76 DW_TAG_inlined_subroutine
.b32 276 // DW_AT_abstract_origin
.b64 $L__tmp112 // DW_AT_low_pc
.b64 $L__tmp114 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 78 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xa28:0x47 DW_TAG_inlined_subroutine
.b32 175 // DW_AT_abstract_origin
.b64 $L__tmp112 // DW_AT_low_pc
.b64 $L__tmp113 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xa3f:0x2f DW_TAG_inlined_subroutine
.b32 156 // DW_AT_abstract_origin
.b64 $L__tmp112 // DW_AT_low_pc
.b64 $L__tmp113 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xa56:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp112 // DW_AT_low_pc
.b64 $L__tmp113 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 7 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0xa6f:0x17 DW_TAG_inlined_subroutine
.b32 214 // DW_AT_abstract_origin
.b64 $L__tmp113 // DW_AT_low_pc
.b64 $L__tmp114 // DW_AT_high_pc
.b8 10 // DW_AT_call_file
.b8 56 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xa89:0x47 DW_TAG_inlined_subroutine
.b32 371 // DW_AT_abstract_origin
.b64 $L__tmp85 // DW_AT_low_pc
.b64 $L__tmp86 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 194 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xaa0:0x2f DW_TAG_inlined_subroutine
.b32 352 // DW_AT_abstract_origin
.b64 $L__tmp85 // DW_AT_low_pc
.b64 $L__tmp86 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xab7:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp85 // DW_AT_low_pc
.b64 $L__tmp86 // DW_AT_high_pc
.b8 13 // DW_AT_call_file
.b8 38 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xad0:0x47 DW_TAG_inlined_subroutine
.b32 321 // DW_AT_abstract_origin
.b64 $L__tmp86 // DW_AT_low_pc
.b64 $L__tmp87 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 195 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xae7:0x2f DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp86 // DW_AT_low_pc
.b64 $L__tmp87 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 81 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xafe:0x17 DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp86 // DW_AT_low_pc
.b64 $L__tmp87 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 180 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 5 // Abbrev [5] 0xb17:0x47 DW_TAG_inlined_subroutine
.b32 321 // DW_AT_abstract_origin
.b64 $L__tmp118 // DW_AT_low_pc
.b64 $L__tmp119 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 199 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xb2e:0x2f DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp118 // DW_AT_low_pc
.b64 $L__tmp119 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 79 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xb45:0x17 DW_TAG_inlined_subroutine
.b32 304 // DW_AT_abstract_origin
.b64 $L__tmp118 // DW_AT_low_pc
.b64 $L__tmp119 // DW_AT_high_pc
.b8 3 // DW_AT_call_file
.b8 180 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 6 // Abbrev [6] 0xb5e:0x17 DW_TAG_inlined_subroutine
.b32 424 // DW_AT_abstract_origin
.b64 $L__tmp119 // DW_AT_low_pc
.b64 $L__tmp120 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 200 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xb75:0x5f DW_TAG_inlined_subroutine
.b32 477 // DW_AT_abstract_origin
.b64 $L__tmp120 // DW_AT_low_pc
.b64 $L__tmp121 // DW_AT_high_pc
.b8 2 // DW_AT_call_file
.b8 203 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xb8c:0x47 DW_TAG_inlined_subroutine
.b32 465 // DW_AT_abstract_origin
.b64 $L__tmp120 // DW_AT_low_pc
.b64 $L__tmp121 // DW_AT_high_pc
.b8 17 // DW_AT_call_file
.b8 8 // DW_AT_call_line
.b8 5 // Abbrev [5] 0xba3:0x2f DW_TAG_inlined_subroutine
.b32 446 // DW_AT_abstract_origin
.b64 $L__tmp120 // DW_AT_low_pc
.b64 $L__tmp121 // DW_AT_high_pc
.b8 18 // DW_AT_call_file
.b8 3 // DW_AT_call_line
.b8 6 // Abbrev [6] 0xbba:0x17 DW_TAG_inlined_subroutine
.b32 84 // DW_AT_abstract_origin
.b64 $L__tmp120 // DW_AT_low_pc
.b64 $L__tmp121 // DW_AT_high_pc
.b8 18 // DW_AT_call_file
.b8 3 // DW_AT_call_line
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b32 243 // Length of Unit
.b8 2 // DWARF version number
.b8 0
.b32 .debug_abbrev // Offset Into Abbrev. Section
.b8 8 // Address Size (in bytes)
.b8 2 // Abbrev [2] 0xb:0xec DW_TAG_compile_unit
.b8 106 // DW_AT_producer
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b8 31 // DW_AT_language
.b8 0
.b8 106 // DW_AT_name
.b8 117
.b8 108
.b8 105
.b8 97
.b8 0
.b32 .debug_line // DW_AT_stmt_list
.b8 46 // DW_AT_comp_dir
.b8 0
.b64 $L__func_begin5 // DW_AT_low_pc
.b64 $L__func_end5 // DW_AT_high_pc
.b8 3 // Abbrev [3] 0x30:0x13 DW_TAG_subprogram
.b8 109 // DW_AT_name
.b8 97
.b8 99
.b8 114
.b8 111
.b8 32
.b8 101
.b8 120
.b8 112
.b8 97
.b8 110
.b8 115
.b8 105
.b8 111
.b8 110
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x43:0xf DW_TAG_subprogram
.b8 101 // DW_AT_name
.b8 110
.b8 122
.b8 121
.b8 109
.b8 101
.b8 95
.b8 99
.b8 97
.b8 108
.b8 108
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x52:0x18 DW_TAG_subprogram
.b8 67 // DW_AT_name
.b8 111
.b8 109
.b8 98
.b8 105
.b8 110
.b8 101
.b8 100
.b8 65
.b8 100
.b8 106
.b8 111
.b8 105
.b8 110
.b8 116
.b8 84
.b8 104
.b8 117
.b8 110
.b8 107
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 3 // Abbrev [3] 0x6a:0x15 DW_TAG_subprogram
.b8 97 // DW_AT_name
.b8 117
.b8 116
.b8 111
.b8 100
.b8 105
.b8 102
.b8 102
.b8 95
.b8 100
.b8 101
.b8 102
.b8 101
.b8 114
.b8 114
.b8 101
.b8 100
.b8 59
.b8 0
.b8 1 // DW_AT_inline
.b8 4 // Abbrev [4] 0x7f:0x77 DW_TAG_subprogram
.b64 $L__func_begin5 // DW_AT_low_pc
.b64 $L__func_end5 // DW_AT_high_pc
.b8 100 // DW_AT_name
.b8 102
.b8 0
.b8 5 // Abbrev [5] 0x93:0x62 DW_TAG_inlined_subroutine
.b32 106 // DW_AT_abstract_origin
.b64 $L__tmp123 // DW_AT_low_pc
.b64 $L__tmp124 // DW_AT_high_pc
.b8 14 // DW_AT_call_file
.b8 14 // DW_AT_call_line
.b8 8 // Abbrev [8] 0xaa:0x4a DW_TAG_inlined_subroutine
.b32 82 // DW_AT_abstract_origin
.b64 $L__tmp123 // DW_AT_low_pc
.b64 $L__tmp124 // DW_AT_high_pc
.b8 20 // DW_AT_call_file
.b8 18 // DW_AT_call_line
.b8 3
.b8 8 // Abbrev [8] 0xc2:0x31 DW_TAG_inlined_subroutine
.b32 67 // DW_AT_abstract_origin
.b64 $L__tmp123 // DW_AT_low_pc
.b64 $L__tmp124 // DW_AT_high_pc
.b8 19 // DW_AT_call_file
.b8 30 // DW_AT_call_line
.b8 19
.b8 7 // Abbrev [7] 0xda:0x18 DW_TAG_inlined_subroutine
.b32 48 // DW_AT_abstract_origin
.b64 $L__tmp123 // DW_AT_low_pc
.b64 $L__tmp124 // DW_AT_high_pc
.b8 19 // DW_AT_call_file
.b8 155 // DW_AT_call_line
.b8 19
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
.b8 0 // End Of Children Mark
}
.section .debug_loc { }
Note for self:
%18 = call {}*** @julia.get_pgcstack(), !dbg !42
%19 = inttoptr i64 ptrtoint (void ({ i8 addrspace(1)*, i64, [1 x i64], i64 } addrspace(11)*)* @julia_f_20562 to i64) to void ({ i8 addrspace(1)*, i64, [1 x i64], i64 }, { i8 addrspace(1)*, i64, [1 x i64], i64 })*, !dbg !42
call void %19({ i8 addrspace(1)*, i64, [1 x i64], i64 } %16, { i8 addrspace(1)*, i64, [1 x i64], i64 } %17), !dbg !42
; └└└└
; @ /home/vchuravy/src/Enzyme/CUDA.jl:15 within `df`
Looks like we are not running Enzyme over the primal code.
Your PR requires formatting changes to meet the project's style guidelines.
Please consider running Runic (git runic main) to apply these changes.
Click here to view the suggested changes.
diff --git a/src/compiler.jl b/src/compiler.jl
index 4a05fbe..e0fd458 100644
--- a/src/compiler.jl
+++ b/src/compiler.jl
@@ -1275,7 +1275,7 @@ end
# Define EnzymeTarget
# Base.@kwdef
-struct EnzymeTarget{Target<:AbstractCompilerTarget} <: AbstractCompilerTarget
+struct EnzymeTarget{Target <: AbstractCompilerTarget} <: AbstractCompilerTarget
target::Target
end
@@ -1292,13 +1292,13 @@ DefaultCompilerTarget(; kwargs...) =
# TODO: Audit uses
function EnzymeTarget()
- EnzymeTarget(DefaultCompilerTarget())
+ return EnzymeTarget(DefaultCompilerTarget())
end
module Runtime end
abstract type AbstractEnzymeCompilerParams <: AbstractCompilerParams end
-struct EnzymeCompilerParams{Params<:AbstractCompilerParams} <: AbstractEnzymeCompilerParams
+struct EnzymeCompilerParams{Params <: AbstractCompilerParams} <: AbstractEnzymeCompilerParams
params::Params
TT::Type{<:Tuple}
@@ -1332,12 +1332,14 @@ struct PrimalCompilerParams <: AbstractEnzymeCompilerParams
mode::API.CDerivativeMode
end
-function EnzymeCompilerParams(TT, mode, width, rt, run_enzyme, abiwrap,
- modifiedBetween, returnPrimal, shadowInit,
- expectedTapeType, ABI,
- err_if_func_written, runtimeActivity, strongZero)
+function EnzymeCompilerParams(
+ TT, mode, width, rt, run_enzyme, abiwrap,
+ modifiedBetween, returnPrimal, shadowInit,
+ expectedTapeType, ABI,
+ err_if_func_written, runtimeActivity, strongZero
+ )
params = PrimalCompilerParams(mode)
- EnzymeCompilerParams(
+ return EnzymeCompilerParams(
params,
TT,
mode,
@@ -1356,9 +1358,9 @@ function EnzymeCompilerParams(TT, mode, width, rt, run_enzyme, abiwrap,
)
end
-# FIXME: Use params.parent in more places where we rely on the behavior of the underlying
+# FIXME: Use params.parent in more places where we rely on the behavior of the underlying
function GPUCompiler.nest_params(params::AbstractEnzymeCompilerParams, parent::AbstractCompilerParams)
- EnzymeCompilerParams(
+ return EnzymeCompilerParams(
parent,
params.TT,
params.mode,
@@ -1381,7 +1383,6 @@ struct UnknownTapeType end
-
## job
# TODO: We shouldn't blanket opt-out
@@ -3511,7 +3512,7 @@ const DumpPreOpt = Ref(false)
function GPUCompiler.compile_unhooked(output::Symbol, job::CompilerJob{<:EnzymeTarget})
@assert output == :llvm
-
+
config = job.config
params = config.params
@@ -3539,7 +3540,7 @@ function GPUCompiler.compile_unhooked(output::Symbol, job::CompilerJob{<:EnzymeT
primal_target = (job.config.target::EnzymeTarget).target
primal_params = (job.config.params::EnzymeCompilerParams).params
if primal_target isa GPUCompiler.NativeCompilerTarget
- @assert primal_params isa PrimalCompilerParams
+ @assert primal_params isa PrimalCompilerParams
else
# XXX: This means mode is not propagated and rules are not applied for GPU code.
end
@@ -4364,15 +4365,15 @@ end
parallel = false
process_module = false
device_module = false
- if primal_target isa GPUCompiler.NativeCompilerTarget
- parallel = Base.Threads.nthreads() > 1
+ if primal_target isa GPUCompiler.NativeCompilerTarget
+ parallel = Base.Threads.nthreads() > 1
else
# All other targets are GPU targets
parallel = true
device_module = true
-
+
if primal_target isa GPUCompiler.GCNCompilerTarget ||
- primal_target isa GPUCompiler.MetalCompilerTarget
+ primal_target isa GPUCompiler.MetalCompilerTarget
process_module = true
end
end
@@ -5583,7 +5584,7 @@ const DumpPostOpt = Ref(false)
# actual compilation
function _thunk(job, postopt::Bool = true)::Tuple{LLVM.Module, Vector{Any}, String, Union{String, Nothing}, Type, String}
- config = CompilerConfig(job.config; optimize=false)
+ config = CompilerConfig(job.config; optimize = false)
job = CompilerJob(job.source, config, job.world)
mod, meta = compile(:llvm, job)
adjointf, augmented_primalf = meta.adjointf, meta.augmented_primalf
diff --git a/src/compiler/reflection.jl b/src/compiler/reflection.jl
index 6b117cd..a767eda 100644
--- a/src/compiler/reflection.jl
+++ b/src/compiler/reflection.jl
@@ -5,7 +5,7 @@ function get_job(
@nospecialize(A),
@nospecialize(types);
run_enzyme::Bool = true,
- optimize::Bool = true,
+ optimize::Bool = true,
mode::API.CDerivativeMode = API.DEM_ReverseModeCombined,
dupClosure::Bool = false,
argwrap::Bool = true,
![github-actions[bot] avatar](https://avatars.githubusercontent.com/in/15368?v=4 "Published by a pub.dev verified publisher"){kind=small}
Well this is new...
ERROR: LoadError: Cannot compile MethodInstance for Enzyme.var"##266".f(::Float64) for world 33088; method is only valid from world 33386 onwards
Stacktrace:
[1] error(s::String)
@ Base ./error.jl:35
[2] compile_method_instance(job::GPUCompiler.CompilerJob)
@ GPUCompiler ~/.cache/julia-buildkite-plugin/depots/260b8bb4-e043-429b-b2f1-bbf9045c057a/packages/GPUCompiler/lB4Vb/src/jlgen.jl:690
[3] macro expansion
@ ~/.cache/julia-buildkite-plugin/depots/260b8bb4-e043-429b-b2f1-bbf9045c057a/packages/Tracy/GcShf/src/tracepoint.jl:158 [inlined]
[4] irgen(job::GPUCompiler.CompilerJob)
@ GPUCompiler ~/.cache/julia-buildkite-plugin/depots/260b8bb4-e043-429b-b2f1-bbf9045c057a/packages/GPUCompiler/lB4Vb/src/irgen.jl:4
Benchmark Results
| main | 42d89764ce5ee6... | main / 42d89764ce5ee6... | |
|---|---|---|---|
| basics/overhead | 4.94 ± 0.92 ns | 4.34 ± 0.01 ns | 1.14 ± 0.21 |
| time_to_load | 1.3 ± 0.014 s | 1.28 ± 0.019 s | 1.02 ± 0.019 |
Benchmark Plots
A plot of the benchmark results has been uploaded as an artifact at https://github.com/EnzymeAD/Enzyme.jl/actions/runs/16007792328/artifacts/3443607038.
Fix nested compilation for GPUCompiler 1.6
Does that mean you should set compat for GPUCompiler to v1.6?
Does that mean you should set compat for GPUCompiler to v1.6?
I did that in #2451 to triple-check that the release is backwards compatible.