//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-29920130
// Cuda compilation tools, release 11.3, V11.3.109
// Based on NVVM 7.0.1
//
.version 7.3
.target sm_52, debug
.address_size 64
.global .align 8 .b8 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_P[48] = {1, 0, 0, 0, 0, 192, 8, 133, 0, 0, 0, 48, 68, 93, 11, 23, 0, 72, 9, 186, 47, 98, 243, 30, 143, 19, 245, 0, 243, 217, 34, 26, 59, 73, 161, 108, 192, 5, 59, 198, 234, 16, 197, 23, 70, 58, 174, 1};
.global .align 8 .b8 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_ONE[48] = {104, 255, 255, 255, 255, 255, 205, 2, 177, 255, 255, 127, 131, 159, 64, 81, 242, 63, 125, 138, 169, 179, 125, 159, 5, 99, 124, 110, 183, 151, 78, 123, 232, 132, 60, 128, 191, 149, 244, 76, 154, 244, 253, 226, 97, 102, 141, 0};
.global .align 8 .b8 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_R2[48] = {34, 205, 0, 148, 108, 104, 134, 183, 177, 49, 4, 176, 170, 252, 41, 3, 109, 180, 214, 98, 17, 241, 165, 34, 172, 195, 125, 130, 3, 125, 223, 191, 249, 11, 121, 65, 240, 146, 126, 131, 136, 75, 145, 30, 203, 252, 109, 0};
.global .align 8 .u64 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_p0 = -8860621160618917889;
.func (.param .b32 func_retval0) _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_gt_384EPKyS1_(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_gt_384EPKyS1__param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_gt_384EPKyS1__param_1
)
{
.reg .pred %p<7>;
.reg .b32 %r<14>;
.reg .b64 %rd<19>;
.loc 2 21 0
$L__func_begin0:
.loc 2 21 0
ld.param.u64 %rd1, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_gt_384EPKyS1__param_0];
ld.param.u64 %rd2, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_gt_384EPKyS1__param_1];
$L__tmp0:
.loc 2 22 16
mov.u32 %r8, 5;
mov.b32 %r1, %r8;
$L__tmp1:
.loc 2 22 5
mov.u32 %r12, %r1;
$L__tmp2:
bra.uni $L__BB0_1;
$L__BB0_1:
mov.u32 %r2, %r12;
$L__tmp3:
setp.ge.s32 %p1, %r2, 0;
not.pred %p2, %p1;
@%p2 bra $L__BB0_9;
bra.uni $L__BB0_2;
$L__BB0_2:
$L__tmp4:
.loc 2 23 9
cvt.s64.s32 %rd3, %r2;
shl.b64 %rd4, %rd3, 3;
add.s64 %rd5, %rd1, %rd4;
ld.u64 %rd6, [%rd5];
cvt.s64.s32 %rd7, %r2;
shl.b64 %rd8, %rd7, 3;
add.s64 %rd9, %rd2, %rd8;
ld.u64 %rd10, [%rd9];
setp.lt.u64 %p3, %rd6, %rd10;
not.pred %p4, %p3;
@%p4 bra $L__BB0_4;
bra.uni $L__BB0_3;
$L__BB0_3:
$L__tmp5:
.loc 2 24 13
mov.u32 %r11, 0;
mov.b32 %r3, %r11;
mov.u32 %r13, %r3;
bra.uni $L__BB0_10;
$L__tmp6:
$L__BB0_4:
.loc 2 25 16
cvt.s64.s32 %rd11, %r2;
shl.b64 %rd12, %rd11, 3;
add.s64 %rd13, %rd1, %rd12;
ld.u64 %rd14, [%rd13];
cvt.s64.s32 %rd15, %r2;
shl.b64 %rd16, %rd15, 3;
add.s64 %rd17, %rd2, %rd16;
ld.u64 %rd18, [%rd17];
setp.gt.u64 %p5, %rd14, %rd18;
not.pred %p6, %p5;
@%p6 bra $L__BB0_6;
bra.uni $L__BB0_5;
$L__BB0_5:
$L__tmp7:
.loc 2 26 13
mov.u32 %r10, 1;
mov.b32 %r4, %r10;
mov.u32 %r13, %r4;
bra.uni $L__BB0_10;
$L__BB0_6:
bra.uni $L__BB0_7;
$L__tmp8:
$L__BB0_7:
.loc 2 22 29
bra.uni $L__BB0_8;
$L__BB0_8:
add.s32 %r5, %r2, -1;
$L__tmp9:
mov.u32 %r12, %r5;
$L__tmp10:
bra.uni $L__BB0_1;
$L__tmp11:
$L__BB0_9:
.loc 2 29 5
mov.u32 %r9, 0;
mov.b32 %r6, %r9;
mov.u32 %r13, %r6;
bra.uni $L__BB0_10;
$L__BB0_10:
mov.u32 %r7, %r13;
st.param.b32 [func_retval0+0], %r7;
ret;
$L__tmp12:
$L__func_end0:
}
.func (.param .b32 func_retval0) _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_ge_384EPKyS1_(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_ge_384EPKyS1__param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_ge_384EPKyS1__param_1
)
{
.reg .pred %p<7>;
.reg .b32 %r<14>;
.reg .b64 %rd<19>;
.loc 3 7 0
$L__func_begin1:
.loc 3 7 0
ld.param.u64 %rd1, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_ge_384EPKyS1__param_0];
ld.param.u64 %rd2, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_ge_384EPKyS1__param_1];
$L__tmp13:
.loc 3 8 16
mov.u32 %r8, 5;
mov.b32 %r1, %r8;
$L__tmp14:
.loc 3 8 5
mov.u32 %r12, %r1;
$L__tmp15:
bra.uni $L__BB1_1;
$L__BB1_1:
mov.u32 %r2, %r12;
$L__tmp16:
setp.ge.s32 %p1, %r2, 0;
not.pred %p2, %p1;
@%p2 bra $L__BB1_9;
bra.uni $L__BB1_2;
$L__BB1_2:
$L__tmp17:
.loc 3 9 9
cvt.s64.s32 %rd3, %r2;
shl.b64 %rd4, %rd3, 3;
add.s64 %rd5, %rd1, %rd4;
ld.u64 %rd6, [%rd5];
cvt.s64.s32 %rd7, %r2;
shl.b64 %rd8, %rd7, 3;
add.s64 %rd9, %rd2, %rd8;
ld.u64 %rd10, [%rd9];
setp.lt.u64 %p3, %rd6, %rd10;
not.pred %p4, %p3;
@%p4 bra $L__BB1_4;
bra.uni $L__BB1_3;
$L__BB1_3:
$L__tmp18:
.loc 3 10 13
mov.u32 %r11, 0;
mov.b32 %r3, %r11;
mov.u32 %r13, %r3;
bra.uni $L__BB1_10;
$L__tmp19:
$L__BB1_4:
.loc 3 11 16
cvt.s64.s32 %rd11, %r2;
shl.b64 %rd12, %rd11, 3;
add.s64 %rd13, %rd1, %rd12;
ld.u64 %rd14, [%rd13];
cvt.s64.s32 %rd15, %r2;
shl.b64 %rd16, %rd15, 3;
add.s64 %rd17, %rd2, %rd16;
ld.u64 %rd18, [%rd17];
setp.gt.u64 %p5, %rd14, %rd18;
not.pred %p6, %p5;
@%p6 bra $L__BB1_6;
bra.uni $L__BB1_5;
$L__BB1_5:
$L__tmp20:
.loc 3 12 13
mov.u32 %r10, 1;
mov.b32 %r4, %r10;
mov.u32 %r13, %r4;
bra.uni $L__BB1_10;
$L__BB1_6:
bra.uni $L__BB1_7;
$L__tmp21:
$L__BB1_7:
.loc 3 8 29
bra.uni $L__BB1_8;
$L__BB1_8:
add.s32 %r5, %r2, -1;
$L__tmp22:
mov.u32 %r12, %r5;
$L__tmp23:
bra.uni $L__BB1_1;
$L__tmp24:
$L__BB1_9:
.loc 3 15 5
mov.u32 %r9, 1;
mov.b32 %r6, %r9;
mov.u32 %r13, %r6;
bra.uni $L__BB1_10;
$L__BB1_10:
mov.u32 %r7, %r13;
st.param.b32 [func_retval0+0], %r7;
ret;
$L__tmp25:
$L__func_end1:
}
.func _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2_(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2__param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2__param_1,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2__param_2
)
{
.reg .b64 %rd<22>;
.loc 3 18 0
$L__func_begin2:
.loc 3 18 0
ld.param.u64 %rd1, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2__param_0];
ld.param.u64 %rd2, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2__param_1];
ld.param.u64 %rd3, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2__param_2];
$L__tmp26:
.loc 3 19 4
ld.u64 %rd10, [%rd2];
ld.u64 %rd11, [%rd2+8];
ld.u64 %rd12, [%rd2+16];
ld.u64 %rd13, [%rd2+24];
ld.u64 %rd14, [%rd2+32];
ld.u64 %rd15, [%rd2+40];
ld.u64 %rd16, [%rd3];
ld.u64 %rd17, [%rd3+8];
ld.u64 %rd18, [%rd3+16];
ld.u64 %rd19, [%rd3+24];
ld.u64 %rd20, [%rd3+32];
ld.u64 %rd21, [%rd3+40];
// begin inline asm
sub.cc.u64 %rd4, %rd10, %rd16;
subc.cc.u64 %rd5, %rd11, %rd17;
subc.cc.u64 %rd6, %rd12, %rd18;
subc.cc.u64 %rd7, %rd13, %rd19;
subc.cc.u64 %rd8, %rd14, %rd20;
subc.u64 %rd9, %rd15, %rd21;
// end inline asm
st.u64 [%rd1], %rd4;
st.u64 [%rd1+8], %rd5;
st.u64 [%rd1+16], %rd6;
st.u64 [%rd1+24], %rd7;
st.u64 [%rd1+32], %rd8;
st.u64 [%rd1+40], %rd9;
.loc 3 46 1
ret;
$L__tmp27:
$L__func_end2:
}
.func _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy_param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy_param_1
)
{
.local .align 8 .b8 __local_depot3[48];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<5>;
.reg .b16 %rs<2>;
.reg .b32 %r<2>;
.reg .b64 %rd<11>;
.loc 3 48 0
$L__func_begin3:
.loc 3 48 0
mov.u64 %SPL, __local_depot3;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd5, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy_param_0];
ld.param.u64 %rd6, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy_param_1];
$L__tmp28:
.loc 3 49 9
{ // callseq 0, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd5;
.param .b64 param1;
st.param.b64 [param1+0], %rd6;
.param .b32 retval0;
call.uni (retval0),
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_ge_384EPKyS1_,
(
param0,
param1
);
ld.param.b32 %r1, [retval0+0];
} // callseq 0
setp.ne.s32 %p1, %r1, 0;
not.pred %p2, %p1;
@%p2 bra $L__BB3_4;
bra.uni $L__BB3_1;
$L__BB3_1:
.loc 3 0 9
add.u64 %rd2, %SP, 0;
$L__tmp29:
.loc 3 51 9
{ // callseq 1, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b64 param1;
st.param.b64 [param1+0], %rd5;
.param .b64 param2;
st.param.b64 [param2+0], %rd6;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2_,
(
param0,
param1,
param2
);
} // callseq 1
mov.u64 %rd7, 0;
.loc 3 52 9
mov.pred %p3, 0;
mov.u64 %rd10, %rd7;
@%p3 bra $L__BB3_3;
bra.uni $L__BB3_2;
$L__BB3_2:
.loc 3 0 9
mov.u64 %rd3, %rd10;
add.s64 %rd8, %rd2, %rd3;
ld.u8 %rs1, [%rd8];
add.s64 %rd9, %rd5, %rd3;
st.u8 [%rd9], %rs1;
add.s64 %rd4, %rd3, 1;
setp.lt.u64 %p4, %rd4, 48;
mov.u64 %rd10, %rd4;
@%p4 bra $L__BB3_2;
bra.uni $L__BB3_3;
$L__BB3_3:
.loc 3 52 9
bra.uni $L__BB3_4;
$L__tmp30:
$L__BB3_4:
.loc 3 54 1
ret;
$L__tmp31:
$L__func_end3:
}
.func _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_1,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_2,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_3
)
{
.reg .pred %p<3>;
.reg .b16 %rs<2>;
.reg .b64 %rd<156>;
.loc 3 59 0
$L__func_begin4:
.loc 3 59 0
ld.param.u64 %rd5, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_0];
ld.param.u64 %rd6, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_1];
ld.param.u64 %rd7, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_2];
ld.param.u64 %rd8, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy_param_3];
$L__tmp32:
.loc 3 61 14
ld.u64 %rd146, [%rd6];
mul.lo.s64 %rd23, %rd146, %rd8;
$L__tmp33:
.loc 3 0 14
mov.u64 %rd145, 0;
.loc 3 63 24
mov.b64 %rd147, %rd145;
$L__tmp34:
.loc 3 65 5
ld.u64 %rd24, [%rd6];
ld.u64 %rd25, [%rd6+8];
ld.u64 %rd26, [%rd6+16];
ld.u64 %rd27, [%rd6+24];
ld.u64 %rd28, [%rd6+32];
ld.u64 %rd29, [%rd6+40];
ld.u64 %rd30, [%rd6+48];
ld.u64 %rd17, [%rd7];
ld.u64 %rd18, [%rd7+8];
ld.u64 %rd19, [%rd7+16];
ld.u64 %rd20, [%rd7+24];
ld.u64 %rd21, [%rd7+32];
ld.u64 %rd22, [%rd7+40];
mov.u64 %rd12, %rd27;
mov.u64 %rd14, %rd29;
mov.u64 %rd9, %rd24;
mov.u64 %rd11, %rd26;
mov.u64 %rd13, %rd28;
mov.u64 %rd15, %rd30;
mov.u64 %rd10, %rd25;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd23, %rd17, %rd9;
madc.hi.cc.u64 c, %rd23, %rd17, 0;
addc.cc.u64 t, %rd10, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd10, %rd23, %rd18, t;
madc.hi.cc.u64 c, %rd23, %rd18, nc;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd23, %rd19, t;
madc.hi.cc.u64 c, %rd23, %rd19, nc;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd23, %rd20, t;
madc.hi.cc.u64 c, %rd23, %rd20, nc;
addc.cc.u64 t, %rd13, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd13, %rd23, %rd21, t;
madc.hi.cc.u64 c, %rd23, %rd21, nc;
addc.cc.u64 t, %rd14, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd14, %rd23, %rd22, t;
madc.hi.cc.u64 c, %rd23, %rd22, nc;
addc.cc.u64 %rd15, %rd15, c;
addc.u64 %rd16, 0, 0;
}
// end inline asm
$L__tmp35:
st.u64 [%rd6], %rd9;
st.u64 [%rd6+8], %rd10;
st.u64 [%rd6+16], %rd11;
st.u64 [%rd6+24], %rd12;
st.u64 [%rd6+32], %rd13;
st.u64 [%rd6+40], %rd14;
st.u64 [%rd6+48], %rd15;
.loc 3 121 5
ld.u64 %rd148, [%rd6+8];
mul.lo.s64 %rd45, %rd148, %rd8;
$L__tmp36:
.loc 3 123 5
ld.u64 %rd46, [%rd6+8];
ld.u64 %rd47, [%rd6+16];
ld.u64 %rd48, [%rd6+24];
ld.u64 %rd49, [%rd6+32];
ld.u64 %rd50, [%rd6+40];
ld.u64 %rd51, [%rd6+48];
ld.u64 %rd52, [%rd6+56];
ld.u64 %rd39, [%rd7];
ld.u64 %rd40, [%rd7+8];
ld.u64 %rd41, [%rd7+16];
ld.u64 %rd42, [%rd7+24];
ld.u64 %rd43, [%rd7+32];
ld.u64 %rd44, [%rd7+40];
mov.u64 %rd31, %rd46;
mov.u64 %rd33, %rd48;
mov.u64 %rd35, %rd50;
mov.u64 %rd37, %rd52;
mov.u64 %rd32, %rd47;
mov.u64 %rd38, %rd16;
mov.u64 %rd34, %rd49;
mov.u64 %rd36, %rd51;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd45, %rd39, %rd31;
madc.hi.cc.u64 c, %rd45, %rd39, 0;
addc.cc.u64 t, %rd32, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd32, %rd45, %rd40, t;
madc.hi.cc.u64 c, %rd45, %rd40, nc;
addc.cc.u64 t, %rd33, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd33, %rd45, %rd41, t;
madc.hi.cc.u64 c, %rd45, %rd41, nc;
addc.cc.u64 t, %rd34, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd34, %rd45, %rd42, t;
madc.hi.cc.u64 c, %rd45, %rd42, nc;
addc.cc.u64 t, %rd35, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd35, %rd45, %rd43, t;
madc.hi.cc.u64 c, %rd45, %rd43, nc;
addc.cc.u64 t, %rd36, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd36, %rd45, %rd44, t;
madc.hi.cc.u64 c, %rd45, %rd44, nc;
addc.cc.u64 c, c, %rd38;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd37, %rd37, c;
addc.u64 %rd38, nc, 0;
}
// end inline asm
$L__tmp37:
st.u64 [%rd6+8], %rd31;
st.u64 [%rd6+16], %rd32;
st.u64 [%rd6+24], %rd33;
st.u64 [%rd6+32], %rd34;
st.u64 [%rd6+40], %rd35;
st.u64 [%rd6+48], %rd36;
st.u64 [%rd6+56], %rd37;
.loc 3 180 5
ld.u64 %rd149, [%rd6+16];
mul.lo.s64 %rd68, %rd149, %rd8;
$L__tmp38:
.loc 3 182 5
ld.u64 %rd69, [%rd6+16];
ld.u64 %rd70, [%rd6+24];
ld.u64 %rd71, [%rd6+32];
ld.u64 %rd72, [%rd6+40];
ld.u64 %rd73, [%rd6+48];
ld.u64 %rd74, [%rd6+56];
ld.u64 %rd75, [%rd6+64];
ld.u64 %rd62, [%rd7];
ld.u64 %rd63, [%rd7+8];
ld.u64 %rd64, [%rd7+16];
ld.u64 %rd65, [%rd7+24];
ld.u64 %rd66, [%rd7+32];
ld.u64 %rd67, [%rd7+40];
mov.u64 %rd57, %rd72;
mov.u64 %rd59, %rd74;
mov.u64 %rd54, %rd69;
mov.u64 %rd61, %rd38;
mov.u64 %rd56, %rd71;
mov.u64 %rd58, %rd73;
mov.u64 %rd60, %rd75;
mov.u64 %rd55, %rd70;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd68, %rd62, %rd54;
madc.hi.cc.u64 c, %rd68, %rd62, 0;
addc.cc.u64 t, %rd55, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd55, %rd68, %rd63, t;
madc.hi.cc.u64 c, %rd68, %rd63, nc;
addc.cc.u64 t, %rd56, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd56, %rd68, %rd64, t;
madc.hi.cc.u64 c, %rd68, %rd64, nc;
addc.cc.u64 t, %rd57, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd57, %rd68, %rd65, t;
madc.hi.cc.u64 c, %rd68, %rd65, nc;
addc.cc.u64 t, %rd58, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd58, %rd68, %rd66, t;
madc.hi.cc.u64 c, %rd68, %rd66, nc;
addc.cc.u64 t, %rd59, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd59, %rd68, %rd67, t;
madc.hi.cc.u64 c, %rd68, %rd67, nc;
addc.cc.u64 c, c, %rd61;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd60, %rd60, c;
addc.u64 %rd61, nc, 0;
}
// end inline asm
$L__tmp39:
st.u64 [%rd6+16], %rd54;
st.u64 [%rd6+24], %rd55;
st.u64 [%rd6+32], %rd56;
st.u64 [%rd6+40], %rd57;
st.u64 [%rd6+48], %rd58;
st.u64 [%rd6+56], %rd59;
st.u64 [%rd6+64], %rd60;
.loc 3 240 5
ld.u64 %rd150, [%rd6+24];
mul.lo.s64 %rd91, %rd150, %rd8;
$L__tmp40:
.loc 3 242 5
ld.u64 %rd92, [%rd6+24];
ld.u64 %rd93, [%rd6+32];
ld.u64 %rd94, [%rd6+40];
ld.u64 %rd95, [%rd6+48];
ld.u64 %rd96, [%rd6+56];
ld.u64 %rd97, [%rd6+64];
ld.u64 %rd98, [%rd6+72];
ld.u64 %rd85, [%rd7];
ld.u64 %rd86, [%rd7+8];
ld.u64 %rd87, [%rd7+16];
ld.u64 %rd88, [%rd7+24];
ld.u64 %rd89, [%rd7+32];
ld.u64 %rd90, [%rd7+40];
mov.u64 %rd83, %rd98;
mov.u64 %rd78, %rd93;
mov.u64 %rd80, %rd95;
mov.u64 %rd82, %rd97;
mov.u64 %rd77, %rd92;
mov.u64 %rd84, %rd61;
mov.u64 %rd79, %rd94;
mov.u64 %rd81, %rd96;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd91, %rd85, %rd77;
madc.hi.cc.u64 c, %rd91, %rd85, 0;
addc.cc.u64 t, %rd78, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd78, %rd91, %rd86, t;
madc.hi.cc.u64 c, %rd91, %rd86, nc;
addc.cc.u64 t, %rd79, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd79, %rd91, %rd87, t;
madc.hi.cc.u64 c, %rd91, %rd87, nc;
addc.cc.u64 t, %rd80, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd80, %rd91, %rd88, t;
madc.hi.cc.u64 c, %rd91, %rd88, nc;
addc.cc.u64 t, %rd81, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd81, %rd91, %rd89, t;
madc.hi.cc.u64 c, %rd91, %rd89, nc;
addc.cc.u64 t, %rd82, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd82, %rd91, %rd90, t;
madc.hi.cc.u64 c, %rd91, %rd90, nc;
addc.cc.u64 c, c, %rd84;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd83, %rd83, c;
addc.u64 %rd84, nc, 0;
}
// end inline asm
$L__tmp41:
st.u64 [%rd6+24], %rd77;
st.u64 [%rd6+32], %rd78;
st.u64 [%rd6+40], %rd79;
st.u64 [%rd6+48], %rd80;
st.u64 [%rd6+56], %rd81;
st.u64 [%rd6+64], %rd82;
st.u64 [%rd6+72], %rd83;
.loc 3 299 5
ld.u64 %rd151, [%rd6+32];
mul.lo.s64 %rd114, %rd151, %rd8;
$L__tmp42:
.loc 3 301 5
ld.u64 %rd115, [%rd6+32];
ld.u64 %rd116, [%rd6+40];
ld.u64 %rd117, [%rd6+48];
ld.u64 %rd118, [%rd6+56];
ld.u64 %rd119, [%rd6+64];
ld.u64 %rd120, [%rd6+72];
ld.u64 %rd121, [%rd6+80];
ld.u64 %rd108, [%rd7];
ld.u64 %rd109, [%rd7+8];
ld.u64 %rd110, [%rd7+16];
ld.u64 %rd111, [%rd7+24];
ld.u64 %rd112, [%rd7+32];
ld.u64 %rd113, [%rd7+40];
mov.u64 %rd102, %rd117;
mov.u64 %rd104, %rd119;
mov.u64 %rd106, %rd121;
mov.u64 %rd101, %rd116;
mov.u64 %rd103, %rd118;
mov.u64 %rd105, %rd120;
mov.u64 %rd100, %rd115;
mov.u64 %rd107, %rd84;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd114, %rd108, %rd100;
madc.hi.cc.u64 c, %rd114, %rd108, 0;
addc.cc.u64 t, %rd101, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd101, %rd114, %rd109, t;
madc.hi.cc.u64 c, %rd114, %rd109, nc;
addc.cc.u64 t, %rd102, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd102, %rd114, %rd110, t;
madc.hi.cc.u64 c, %rd114, %rd110, nc;
addc.cc.u64 t, %rd103, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd103, %rd114, %rd111, t;
madc.hi.cc.u64 c, %rd114, %rd111, nc;
addc.cc.u64 t, %rd104, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd104, %rd114, %rd112, t;
madc.hi.cc.u64 c, %rd114, %rd112, nc;
addc.cc.u64 t, %rd105, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd105, %rd114, %rd113, t;
madc.hi.cc.u64 c, %rd114, %rd113, nc;
addc.cc.u64 c, c, %rd107;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd106, %rd106, c;
addc.u64 %rd107, nc, 0;
}
// end inline asm
$L__tmp43:
st.u64 [%rd6+32], %rd100;
st.u64 [%rd6+40], %rd101;
st.u64 [%rd6+48], %rd102;
st.u64 [%rd6+56], %rd103;
st.u64 [%rd6+64], %rd104;
st.u64 [%rd6+72], %rd105;
st.u64 [%rd6+80], %rd106;
.loc 3 358 5
ld.u64 %rd152, [%rd6+40];
mul.lo.s64 %rd137, %rd152, %rd8;
$L__tmp44:
.loc 3 360 5
ld.u64 %rd138, [%rd6+40];
ld.u64 %rd139, [%rd6+48];
ld.u64 %rd140, [%rd6+56];
ld.u64 %rd141, [%rd6+64];
ld.u64 %rd142, [%rd6+72];
ld.u64 %rd143, [%rd6+80];
ld.u64 %rd144, [%rd6+88];
ld.u64 %rd131, [%rd7];
ld.u64 %rd132, [%rd7+8];
ld.u64 %rd133, [%rd7+16];
ld.u64 %rd134, [%rd7+24];
ld.u64 %rd135, [%rd7+32];
ld.u64 %rd136, [%rd7+40];
mov.u64 %rd128, %rd143;
mov.u64 %rd123, %rd138;
mov.u64 %rd125, %rd140;
mov.u64 %rd127, %rd142;
mov.u64 %rd129, %rd144;
mov.u64 %rd124, %rd139;
mov.u64 %rd126, %rd141;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd137, %rd131, %rd123;
madc.hi.cc.u64 c, %rd137, %rd131, 0;
addc.cc.u64 t, %rd124, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd124, %rd137, %rd132, t;
madc.hi.cc.u64 c, %rd137, %rd132, nc;
addc.cc.u64 t, %rd125, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd125, %rd137, %rd133, t;
madc.hi.cc.u64 c, %rd137, %rd133, nc;
addc.cc.u64 t, %rd126, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd126, %rd137, %rd134, t;
madc.hi.cc.u64 c, %rd137, %rd134, nc;
addc.cc.u64 t, %rd127, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd127, %rd137, %rd135, t;
madc.hi.cc.u64 c, %rd137, %rd135, nc;
addc.cc.u64 t, %rd128, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd128, %rd137, %rd136, t;
madc.hi.cc.u64 c, %rd137, %rd136, nc;
addc.cc.u64 c, c, %rd107;
add.u64 %rd129, %rd129, c;
}
// end inline asm
st.u64 [%rd6+40], %rd123;
st.u64 [%rd6+48], %rd124;
st.u64 [%rd6+56], %rd125;
st.u64 [%rd6+64], %rd126;
st.u64 [%rd6+72], %rd127;
st.u64 [%rd6+80], %rd128;
st.u64 [%rd6+88], %rd129;
.loc 3 418 5
add.s64 %rd2, %rd6, 48;
mov.pred %p1, 0;
mov.u64 %rd155, %rd145;
@%p1 bra $L__BB4_2;
bra.uni $L__BB4_1;
$L__BB4_1:
.loc 3 0 5
mov.u64 %rd3, %rd155;
add.s64 %rd153, %rd2, %rd3;
ld.u8 %rs1, [%rd153];
add.s64 %rd154, %rd5, %rd3;
st.u8 [%rd154], %rs1;
add.s64 %rd4, %rd3, 1;
setp.lt.u64 %p2, %rd4, 48;
mov.u64 %rd155, %rd4;
@%p2 bra $L__BB4_1;
bra.uni $L__BB4_2;
$L__BB4_2:
.loc 3 419 5
{ // callseq 2, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd5;
.param .b64 param1;
st.param.b64 [param1+0], %rd7;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy,
(
param0,
param1
);
} // callseq 2
.loc 3 420 1
ret;
$L__tmp45:
$L__func_end4:
}
// .globl _Z12mul_mont_384PyPKyS1_S1_y
.visible .func _Z12mul_mont_384PyPKyS1_S1_y(
.param .b64 _Z12mul_mont_384PyPKyS1_S1_y_param_0,
.param .b64 _Z12mul_mont_384PyPKyS1_S1_y_param_1,
.param .b64 _Z12mul_mont_384PyPKyS1_S1_y_param_2,
.param .b64 _Z12mul_mont_384PyPKyS1_S1_y_param_3,
.param .b64 _Z12mul_mont_384PyPKyS1_S1_y_param_4
)
{
.local .align 8 .b8 __local_depot5[96];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .b64 %rd<43>;
.loc 3 422 0
$L__func_begin5:
.loc 3 422 0
mov.u64 %SPL, __local_depot5;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd1, [_Z12mul_mont_384PyPKyS1_S1_y_param_0];
ld.param.u64 %rd2, [_Z12mul_mont_384PyPKyS1_S1_y_param_1];
ld.param.u64 %rd3, [_Z12mul_mont_384PyPKyS1_S1_y_param_2];
ld.param.u64 %rd4, [_Z12mul_mont_384PyPKyS1_S1_y_param_3];
ld.param.u64 %rd5, [_Z12mul_mont_384PyPKyS1_S1_y_param_4];
$L__tmp46:
.loc 3 425 5
ld.u64 %rd30, [%SP+0];
ld.u64 %rd31, [%SP+8];
ld.u64 %rd32, [%SP+16];
ld.u64 %rd33, [%SP+24];
ld.u64 %rd34, [%SP+32];
ld.u64 %rd35, [%SP+40];
ld.u64 %rd36, [%SP+48];
ld.u64 %rd37, [%SP+56];
ld.u64 %rd38, [%SP+64];
ld.u64 %rd39, [%SP+72];
ld.u64 %rd40, [%SP+80];
ld.u64 %rd41, [%SP+88];
ld.u64 %rd18, [%rd2];
ld.u64 %rd19, [%rd2+8];
ld.u64 %rd20, [%rd2+16];
ld.u64 %rd21, [%rd2+24];
ld.u64 %rd22, [%rd2+32];
ld.u64 %rd23, [%rd2+40];
ld.u64 %rd24, [%rd3];
ld.u64 %rd25, [%rd3+8];
ld.u64 %rd26, [%rd3+16];
ld.u64 %rd27, [%rd3+24];
ld.u64 %rd28, [%rd3+32];
ld.u64 %rd29, [%rd3+40];
mov.u64 %rd10, %rd34;
mov.u64 %rd17, %rd41;
mov.u64 %rd12, %rd36;
mov.u64 %rd7, %rd31;
mov.u64 %rd14, %rd38;
mov.u64 %rd9, %rd33;
mov.u64 %rd16, %rd40;
mov.u64 %rd11, %rd35;
mov.u64 %rd6, %rd30;
mov.u64 %rd13, %rd37;
mov.u64 %rd8, %rd32;
mov.u64 %rd15, %rd39;
// begin inline asm
{
.reg .u64 c;
.reg .u64 nc;
.reg .u64 t;
mad.lo.cc.u64 %rd6, %rd18, %rd24, 0;
madc.hi.cc.u64 c, %rd18, %rd24, 0;
madc.lo.cc.u64 %rd7, %rd18, %rd25, c;
madc.hi.cc.u64 c, %rd18, %rd25, 0;
madc.lo.cc.u64 %rd8, %rd18, %rd26, c;
madc.hi.cc.u64 c, %rd18, %rd26, 0;
madc.lo.cc.u64 %rd9, %rd18, %rd27, c;
madc.hi.cc.u64 c, %rd18, %rd27, 0;
madc.lo.cc.u64 %rd10, %rd18, %rd28, c;
madc.hi.cc.u64 c, %rd18, %rd28, 0;
madc.lo.cc.u64 %rd11, %rd18, %rd29, c;
madc.hi.u64 %rd12, %rd18, %rd29, 0;
mad.lo.cc.u64 %rd7, %rd19, %rd24, %rd7;
madc.hi.cc.u64 c, %rd19, %rd24, 0;
addc.cc.u64 t, %rd8, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd8, %rd19, %rd25, t;
madc.hi.cc.u64 c, %rd19, %rd25, nc;
addc.cc.u64 t, %rd9, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd9, %rd19, %rd26, t;
madc.hi.cc.u64 c, %rd19, %rd26, nc;
addc.cc.u64 t, %rd10, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd10, %rd19, %rd27, t;
madc.hi.cc.u64 c, %rd19, %rd27, nc;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd19, %rd28, t;
madc.hi.cc.u64 c, %rd19, %rd28, nc;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd19, %rd29, t;
madc.hi.u64 %rd13, %rd19, %rd29, nc;
mad.lo.cc.u64 %rd8, %rd20, %rd24, %rd8;
madc.hi.cc.u64 c, %rd20, %rd24, 0;
addc.cc.u64 t, %rd9, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd9, %rd20, %rd25, t;
madc.hi.cc.u64 c, %rd20, %rd25, nc;
addc.cc.u64 t, %rd10, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd10, %rd20, %rd26, t;
madc.hi.cc.u64 c, %rd20, %rd26, nc;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd20, %rd27, t;
madc.hi.cc.u64 c, %rd20, %rd27, nc;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd20, %rd28, t;
madc.hi.cc.u64 c, %rd20, %rd28, nc;
addc.cc.u64 t, %rd13, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd13, %rd20, %rd29, t;
madc.hi.u64 %rd14, %rd20, %rd29, nc;
mad.lo.cc.u64 %rd9, %rd21, %rd24, %rd9;
madc.hi.cc.u64 c, %rd21, %rd24, 0;
addc.cc.u64 t, %rd10, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd10, %rd21, %rd25, t;
madc.hi.cc.u64 c, %rd21, %rd25, nc;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd21, %rd26, t;
madc.hi.cc.u64 c, %rd21, %rd26, nc;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd21, %rd27, t;
madc.hi.cc.u64 c, %rd21, %rd27, nc;
addc.cc.u64 t, %rd13, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd13, %rd21, %rd28, t;
madc.hi.cc.u64 c, %rd21, %rd28, nc;
addc.cc.u64 t, %rd14, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd14, %rd21, %rd29, t;
madc.hi.u64 %rd15, %rd21, %rd29, nc;
mad.lo.cc.u64 %rd10, %rd22, %rd24, %rd10;
madc.hi.cc.u64 c, %rd22, %rd24, 0;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd22, %rd25, t;
madc.hi.cc.u64 c, %rd22, %rd25, nc;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd22, %rd26, t;
madc.hi.cc.u64 c, %rd22, %rd26, nc;
addc.cc.u64 t, %rd13, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd13, %rd22, %rd27, t;
madc.hi.cc.u64 c, %rd22, %rd27, nc;
addc.cc.u64 t, %rd14, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd14, %rd22, %rd28, t;
madc.hi.cc.u64 c, %rd22, %rd28, nc;
addc.cc.u64 t, %rd15, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd15, %rd22, %rd29, t;
madc.hi.u64 %rd16, %rd22, %rd29, nc;
mad.lo.cc.u64 %rd11, %rd23, %rd24, %rd11;
madc.hi.cc.u64 c, %rd23, %rd24, 0;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd23, %rd25, t;
madc.hi.cc.u64 c, %rd23, %rd25, nc;
addc.cc.u64 t, %rd13, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd13, %rd23, %rd26, t;
madc.hi.cc.u64 c, %rd23, %rd26, nc;
addc.cc.u64 t, %rd14, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd14, %rd23, %rd27, t;
madc.hi.cc.u64 c, %rd23, %rd27, nc;
addc.cc.u64 t, %rd15, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd15, %rd23, %rd28, t;
madc.hi.cc.u64 c, %rd23, %rd28, nc;
addc.cc.u64 t, %rd16, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd16, %rd23, %rd29, t;
madc.hi.u64 %rd17, %rd23, %rd29, nc;
}
// end inline asm
st.u64 [%SP+0], %rd6;
st.u64 [%SP+8], %rd7;
st.u64 [%SP+16], %rd8;
st.u64 [%SP+24], %rd9;
st.u64 [%SP+32], %rd10;
st.u64 [%SP+40], %rd11;
st.u64 [%SP+48], %rd12;
st.u64 [%SP+56], %rd13;
st.u64 [%SP+64], %rd14;
st.u64 [%SP+72], %rd15;
st.u64 [%SP+80], %rd16;
st.u64 [%SP+88], %rd17;
add.u64 %rd42, %SP, 0;
.loc 3 625 5
{ // callseq 3, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd42;
.param .b64 param2;
st.param.b64 [param2+0], %rd4;
.param .b64 param3;
st.param.b64 [param3+0], %rd5;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy,
(
param0,
param1,
param2,
param3
);
} // callseq 3
.loc 3 626 1
ret;
$L__tmp47:
$L__func_end5:
}
// .globl _Z12sqr_mont_384PyPKyS1_y
.visible .func _Z12sqr_mont_384PyPKyS1_y(
.param .b64 _Z12sqr_mont_384PyPKyS1_y_param_0,
.param .b64 _Z12sqr_mont_384PyPKyS1_y_param_1,
.param .b64 _Z12sqr_mont_384PyPKyS1_y_param_2,
.param .b64 _Z12sqr_mont_384PyPKyS1_y_param_3
)
{
.local .align 8 .b8 __local_depot6[96];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .b64 %rd<115>;
.loc 3 628 0
$L__func_begin6:
.loc 3 628 0
mov.u64 %SPL, __local_depot6;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd1, [_Z12sqr_mont_384PyPKyS1_y_param_0];
ld.param.u64 %rd2, [_Z12sqr_mont_384PyPKyS1_y_param_1];
ld.param.u64 %rd3, [_Z12sqr_mont_384PyPKyS1_y_param_2];
ld.param.u64 %rd4, [_Z12sqr_mont_384PyPKyS1_y_param_3];
$L__tmp48:
.loc 3 631 5
ld.u64 %rd23, [%SP+0];
ld.u64 %rd24, [%SP+8];
ld.u64 %rd25, [%SP+16];
ld.u64 %rd26, [%SP+24];
ld.u64 %rd27, [%SP+32];
ld.u64 %rd28, [%SP+40];
ld.u64 %rd29, [%SP+48];
ld.u64 %rd30, [%SP+56];
ld.u64 %rd31, [%SP+64];
ld.u64 %rd32, [%SP+72];
ld.u64 %rd33, [%SP+80];
ld.u64 %rd34, [%SP+88];
ld.u64 %rd17, [%rd2];
ld.u64 %rd18, [%rd2+8];
ld.u64 %rd19, [%rd2+16];
ld.u64 %rd20, [%rd2+24];
ld.u64 %rd21, [%rd2+32];
ld.u64 %rd22, [%rd2+40];
mov.u64 %rd9, %rd27;
mov.u64 %rd16, %rd34;
mov.u64 %rd11, %rd29;
mov.u64 %rd6, %rd24;
mov.u64 %rd13, %rd31;
mov.u64 %rd8, %rd26;
mov.u64 %rd15, %rd33;
mov.u64 %rd10, %rd28;
mov.u64 %rd5, %rd23;
mov.u64 %rd12, %rd30;
mov.u64 %rd7, %rd25;
mov.u64 %rd14, %rd32;
// begin inline asm
{
.reg .u64 c;
.reg .u64 nc;
.reg .u64 t;
mad.lo.cc.u64 %rd6, %rd17, %rd18, 0;
madc.hi.cc.u64 c, %rd17, %rd18, 0;
madc.lo.cc.u64 %rd7, %rd17, %rd19, c;
madc.hi.cc.u64 c, %rd17, %rd19, 0;
madc.lo.cc.u64 %rd8, %rd17, %rd20, c;
madc.hi.cc.u64 c, %rd17, %rd20, 0;
madc.lo.cc.u64 %rd9, %rd17, %rd21, c;
madc.hi.cc.u64 c, %rd17, %rd21, 0;
madc.lo.cc.u64 %rd10, %rd17, %rd22, c;
madc.hi.u64 %rd11, %rd17, %rd22, 0;
mad.lo.cc.u64 %rd8, %rd18, %rd19, %rd8;
madc.hi.cc.u64 c, %rd18, %rd19, 0;
addc.cc.u64 t, %rd9, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd9, %rd18, %rd20, t;
madc.hi.cc.u64 c, %rd18, %rd20, nc;
addc.cc.u64 t, %rd10, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd10, %rd18, %rd21, t;
madc.hi.cc.u64 c, %rd18, %rd21, nc;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd18, %rd22, t;
madc.hi.u64 %rd12, %rd18, %rd22, nc;
mad.lo.cc.u64 %rd10, %rd19, %rd20, %rd10;
madc.hi.cc.u64 c, %rd19, %rd20, 0;
addc.cc.u64 t, %rd11, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd11, %rd19, %rd21, t;
madc.hi.cc.u64 c, %rd19, %rd21, nc;
addc.cc.u64 t, %rd12, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd12, %rd19, %rd22, t;
madc.hi.u64 %rd13, %rd19, %rd22, nc;
mad.lo.cc.u64 %rd12, %rd20, %rd21, %rd12;
madc.hi.cc.u64 c, %rd20, %rd21, 0;
addc.cc.u64 t, %rd13, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd13, %rd20, %rd22, t;
madc.hi.u64 %rd14, %rd20, %rd22, nc;
mad.lo.cc.u64 %rd14, %rd21, %rd22, %rd14;
madc.hi.u64 %rd15, %rd21, %rd22, 0;
}
// end inline asm
st.u64 [%SP+0], %rd5;
st.u64 [%SP+8], %rd6;
st.u64 [%SP+16], %rd7;
st.u64 [%SP+24], %rd8;
st.u64 [%SP+32], %rd9;
st.u64 [%SP+40], %rd10;
st.u64 [%SP+48], %rd11;
st.u64 [%SP+56], %rd12;
st.u64 [%SP+64], %rd13;
st.u64 [%SP+72], %rd14;
st.u64 [%SP+80], %rd15;
st.u64 [%SP+88], %rd16;
.loc 3 724 5
ld.u64 %rd65, [%SP+80];
shr.u64 %rd66, %rd65, 63;
st.u64 [%SP+88], %rd66;
.loc 3 725 5
ld.u64 %rd67, [%SP+80];
shl.b64 %rd68, %rd67, 1;
ld.u64 %rd69, [%SP+72];
shr.u64 %rd70, %rd69, 63;
or.b64 %rd71, %rd68, %rd70;
st.u64 [%SP+80], %rd71;
.loc 3 726 5
ld.u64 %rd72, [%SP+72];
shl.b64 %rd73, %rd72, 1;
ld.u64 %rd74, [%SP+64];
shr.u64 %rd75, %rd74, 63;
or.b64 %rd76, %rd73, %rd75;
st.u64 [%SP+72], %rd76;
.loc 3 727 5
ld.u64 %rd77, [%SP+64];
shl.b64 %rd78, %rd77, 1;
ld.u64 %rd79, [%SP+56];
shr.u64 %rd80, %rd79, 63;
or.b64 %rd81, %rd78, %rd80;
st.u64 [%SP+64], %rd81;
.loc 3 728 5
ld.u64 %rd82, [%SP+56];
shl.b64 %rd83, %rd82, 1;
ld.u64 %rd84, [%SP+48];
shr.u64 %rd85, %rd84, 63;
or.b64 %rd86, %rd83, %rd85;
st.u64 [%SP+56], %rd86;
.loc 3 729 5
ld.u64 %rd87, [%SP+48];
shl.b64 %rd88, %rd87, 1;
ld.u64 %rd89, [%SP+40];
shr.u64 %rd90, %rd89, 63;
or.b64 %rd91, %rd88, %rd90;
st.u64 [%SP+48], %rd91;
.loc 3 730 5
ld.u64 %rd92, [%SP+40];
shl.b64 %rd93, %rd92, 1;
ld.u64 %rd94, [%SP+32];
shr.u64 %rd95, %rd94, 63;
or.b64 %rd96, %rd93, %rd95;
st.u64 [%SP+40], %rd96;
.loc 3 731 5
ld.u64 %rd97, [%SP+32];
shl.b64 %rd98, %rd97, 1;
ld.u64 %rd99, [%SP+24];
shr.u64 %rd100, %rd99, 63;
or.b64 %rd101, %rd98, %rd100;
st.u64 [%SP+32], %rd101;
.loc 3 732 5
ld.u64 %rd102, [%SP+24];
shl.b64 %rd103, %rd102, 1;
ld.u64 %rd104, [%SP+16];
shr.u64 %rd105, %rd104, 63;
or.b64 %rd106, %rd103, %rd105;
st.u64 [%SP+24], %rd106;
.loc 3 733 5
ld.u64 %rd107, [%SP+16];
shl.b64 %rd108, %rd107, 1;
ld.u64 %rd109, [%SP+8];
shr.u64 %rd110, %rd109, 63;
or.b64 %rd111, %rd108, %rd110;
st.u64 [%SP+16], %rd111;
.loc 3 734 5
ld.u64 %rd112, [%SP+8];
shl.b64 %rd113, %rd112, 1;
st.u64 [%SP+8], %rd113;
.loc 3 738 5
ld.u64 %rd53, [%SP+0];
ld.u64 %rd54, [%SP+8];
ld.u64 %rd55, [%SP+16];
ld.u64 %rd56, [%SP+24];
ld.u64 %rd57, [%SP+32];
ld.u64 %rd58, [%SP+40];
ld.u64 %rd59, [%SP+48];
ld.u64 %rd60, [%SP+56];
ld.u64 %rd61, [%SP+64];
ld.u64 %rd62, [%SP+72];
ld.u64 %rd63, [%SP+80];
ld.u64 %rd64, [%SP+88];
ld.u64 %rd47, [%rd2];
ld.u64 %rd48, [%rd2+8];
ld.u64 %rd49, [%rd2+16];
ld.u64 %rd50, [%rd2+24];
ld.u64 %rd51, [%rd2+32];
ld.u64 %rd52, [%rd2+40];
mov.u64 %rd35, %rd53;
mov.u64 %rd42, %rd60;
mov.u64 %rd37, %rd55;
mov.u64 %rd44, %rd62;
mov.u64 %rd39, %rd57;
mov.u64 %rd46, %rd64;
mov.u64 %rd41, %rd59;
mov.u64 %rd36, %rd54;
mov.u64 %rd43, %rd61;
mov.u64 %rd38, %rd56;
mov.u64 %rd45, %rd63;
mov.u64 %rd40, %rd58;
// begin inline asm
{
mad.lo.cc.u64 %rd35, %rd47, %rd47, 0;
madc.hi.cc.u64 %rd36, %rd47, %rd47, %rd36;
madc.lo.cc.u64 %rd37, %rd48, %rd48, %rd37;
madc.hi.cc.u64 %rd38, %rd48, %rd48, %rd38;
madc.lo.cc.u64 %rd39, %rd49, %rd49, %rd39;
madc.hi.cc.u64 %rd40, %rd49, %rd49, %rd40;
madc.lo.cc.u64 %rd41, %rd50, %rd50, %rd41;
madc.hi.cc.u64 %rd42, %rd50, %rd50, %rd42;
madc.lo.cc.u64 %rd43, %rd51, %rd51, %rd43;
madc.hi.cc.u64 %rd44, %rd51, %rd51, %rd44;
madc.lo.cc.u64 %rd45, %rd52, %rd52, %rd45;
madc.hi.u64 %rd46, %rd52, %rd52, %rd46;
}
// end inline asm
st.u64 [%SP+0], %rd35;
st.u64 [%SP+8], %rd36;
st.u64 [%SP+16], %rd37;
st.u64 [%SP+24], %rd38;
st.u64 [%SP+32], %rd39;
st.u64 [%SP+40], %rd40;
st.u64 [%SP+48], %rd41;
st.u64 [%SP+56], %rd42;
st.u64 [%SP+64], %rd43;
st.u64 [%SP+72], %rd44;
st.u64 [%SP+80], %rd45;
st.u64 [%SP+88], %rd46;
add.u64 %rd114, %SP, 0;
.loc 3 781 5
{ // callseq 4, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd114;
.param .b64 param2;
st.param.b64 [param2+0], %rd3;
.param .b64 param3;
st.param.b64 [param3+0], %rd4;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf8mont_384EPyS0_PKyy,
(
param0,
param1,
param2,
param3
);
} // callseq 4
.loc 3 782 1
ret;
$L__tmp49:
$L__func_end6:
}
.func _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2_(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2__param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2__param_1,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2__param_2
)
{
.reg .b64 %rd<22>;
.loc 3 785 0
$L__func_begin7:
.loc 3 785 0
ld.param.u64 %rd1, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2__param_0];
ld.param.u64 %rd2, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2__param_1];
ld.param.u64 %rd3, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2__param_2];
$L__tmp50:
.loc 3 786 5
ld.u64 %rd10, [%rd2];
ld.u64 %rd11, [%rd2+8];
ld.u64 %rd12, [%rd2+16];
ld.u64 %rd13, [%rd2+24];
ld.u64 %rd14, [%rd2+32];
ld.u64 %rd15, [%rd2+40];
ld.u64 %rd16, [%rd3];
ld.u64 %rd17, [%rd3+8];
ld.u64 %rd18, [%rd3+16];
ld.u64 %rd19, [%rd3+24];
ld.u64 %rd20, [%rd3+32];
ld.u64 %rd21, [%rd3+40];
// begin inline asm
add.cc.u64 %rd4, %rd10, %rd16;
addc.cc.u64 %rd5, %rd11, %rd17;
addc.cc.u64 %rd6, %rd12, %rd18;
addc.cc.u64 %rd7, %rd13, %rd19;
addc.cc.u64 %rd8, %rd14, %rd20;
addc.u64 %rd9, %rd15, %rd21;
// end inline asm
st.u64 [%rd1], %rd4;
st.u64 [%rd1+8], %rd5;
st.u64 [%rd1+16], %rd6;
st.u64 [%rd1+24], %rd7;
st.u64 [%rd1+32], %rd8;
st.u64 [%rd1+40], %rd9;
.loc 3 813 1
ret;
$L__tmp51:
$L__func_end7:
}
// .globl _Z11add_mod_384PyPKyS1_S1_
.visible .func _Z11add_mod_384PyPKyS1_S1_(
.param .b64 _Z11add_mod_384PyPKyS1_S1__param_0,
.param .b64 _Z11add_mod_384PyPKyS1_S1__param_1,
.param .b64 _Z11add_mod_384PyPKyS1_S1__param_2,
.param .b64 _Z11add_mod_384PyPKyS1_S1__param_3
)
{
.reg .b64 %rd<5>;
.loc 3 815 0
$L__func_begin8:
.loc 3 815 0
ld.param.u64 %rd1, [_Z11add_mod_384PyPKyS1_S1__param_0];
ld.param.u64 %rd2, [_Z11add_mod_384PyPKyS1_S1__param_1];
ld.param.u64 %rd3, [_Z11add_mod_384PyPKyS1_S1__param_2];
ld.param.u64 %rd4, [_Z11add_mod_384PyPKyS1_S1__param_3];
$L__tmp52:
.loc 3 816 5
{ // callseq 5, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd2;
.param .b64 param2;
st.param.b64 [param2+0], %rd3;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2_,
(
param0,
param1,
param2
);
} // callseq 5
.loc 3 818 5
{ // callseq 6, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd4;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf6reduceEPyPKy,
(
param0,
param1
);
} // callseq 6
.loc 3 820 1
ret;
$L__tmp53:
$L__func_end8:
}
// .globl _Z11sub_mod_384PyPKyS1_S1_
.visible .func _Z11sub_mod_384PyPKyS1_S1_(
.param .b64 _Z11sub_mod_384PyPKyS1_S1__param_0,
.param .b64 _Z11sub_mod_384PyPKyS1_S1__param_1,
.param .b64 _Z11sub_mod_384PyPKyS1_S1__param_2,
.param .b64 _Z11sub_mod_384PyPKyS1_S1__param_3
)
{
.local .align 8 .b8 __local_depot9[48];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<5>;
.reg .b16 %rs<2>;
.reg .b32 %r<2>;
.reg .b64 %rd<15>;
.loc 3 822 0
$L__func_begin9:
.loc 3 822 0
mov.u64 %SPL, __local_depot9;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd5, [_Z11sub_mod_384PyPKyS1_S1__param_0];
ld.param.u64 %rd6, [_Z11sub_mod_384PyPKyS1_S1__param_1];
ld.param.u64 %rd7, [_Z11sub_mod_384PyPKyS1_S1__param_2];
ld.param.u64 %rd8, [_Z11sub_mod_384PyPKyS1_S1__param_3];
add.u64 %rd1, %SP, 0;
mov.u64 %rd9, 0;
$L__tmp54:
.loc 3 824 5
mov.pred %p1, 0;
mov.u64 %rd14, %rd9;
@%p1 bra $L__BB9_2;
bra.uni $L__BB9_1;
$L__BB9_1:
.loc 3 0 5
mov.u64 %rd3, %rd14;
add.s64 %rd10, %rd6, %rd3;
ld.u8 %rs1, [%rd10];
add.s64 %rd11, %rd1, %rd3;
st.u8 [%rd11], %rs1;
add.s64 %rd4, %rd3, 1;
setp.lt.u64 %p2, %rd4, 48;
mov.u64 %rd14, %rd4;
@%p2 bra $L__BB9_1;
bra.uni $L__BB9_2;
$L__BB9_2:
.loc 3 826 9
{ // callseq 7, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd7;
.param .b64 param1;
st.param.b64 [param1+0], %rd6;
.param .b32 retval0;
call.uni (retval0),
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf9is_gt_384EPKyS1_,
(
param0,
param1
);
ld.param.b32 %r1, [retval0+0];
} // callseq 7
setp.ne.s32 %p3, %r1, 0;
not.pred %p4, %p3;
@%p4 bra $L__BB9_4;
bra.uni $L__BB9_3;
$L__BB9_3:
.loc 3 0 9
add.u64 %rd12, %SP, 0;
$L__tmp55:
.loc 3 828 7
{ // callseq 8, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd12;
.param .b64 param1;
st.param.b64 [param1+0], %rd12;
.param .b64 param2;
st.param.b64 [param2+0], %rd8;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2_,
(
param0,
param1,
param2
);
} // callseq 8
bra.uni $L__BB9_5;
$L__tmp56:
$L__BB9_4:
.loc 3 830 12
bra.uni $L__BB9_5;
$L__BB9_5:
.loc 3 0 12
add.u64 %rd13, %SP, 0;
.loc 3 833 5
{ // callseq 9, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd5;
.param .b64 param1;
st.param.b64 [param1+0], %rd13;
.param .b64 param2;
st.param.b64 [param2+0], %rd7;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2_,
(
param0,
param1,
param2
);
} // callseq 9
.loc 3 836 1
ret;
$L__tmp57:
$L__func_end9:
}
// .globl _Z18sub_mod_384_unsafePyPKyS1_
.visible .func _Z18sub_mod_384_unsafePyPKyS1_(
.param .b64 _Z18sub_mod_384_unsafePyPKyS1__param_0,
.param .b64 _Z18sub_mod_384_unsafePyPKyS1__param_1,
.param .b64 _Z18sub_mod_384_unsafePyPKyS1__param_2
)
{
.reg .b64 %rd<4>;
.loc 3 838 0
$L__func_begin10:
.loc 3 838 0
ld.param.u64 %rd1, [_Z18sub_mod_384_unsafePyPKyS1__param_0];
ld.param.u64 %rd2, [_Z18sub_mod_384_unsafePyPKyS1__param_1];
ld.param.u64 %rd3, [_Z18sub_mod_384_unsafePyPKyS1__param_2];
$L__tmp58:
.loc 3 839 5
{ // callseq 10, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd2;
.param .b64 param2;
st.param.b64 [param2+0], %rd3;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21sub_mod_384_uncheckedEPyPKyS2_,
(
param0,
param1,
param2
);
} // callseq 10
.loc 3 841 1
ret;
$L__tmp59:
$L__func_end10:
}
// .globl _Z18add_mod_384_unsafePyPKyS1_
.visible .func _Z18add_mod_384_unsafePyPKyS1_(
.param .b64 _Z18add_mod_384_unsafePyPKyS1__param_0,
.param .b64 _Z18add_mod_384_unsafePyPKyS1__param_1,
.param .b64 _Z18add_mod_384_unsafePyPKyS1__param_2
)
{
.reg .b64 %rd<4>;
.loc 3 843 0
$L__func_begin11:
.loc 3 843 0
ld.param.u64 %rd1, [_Z18add_mod_384_unsafePyPKyS1__param_0];
ld.param.u64 %rd2, [_Z18add_mod_384_unsafePyPKyS1__param_1];
ld.param.u64 %rd3, [_Z18add_mod_384_unsafePyPKyS1__param_2];
$L__tmp60:
.loc 3 844 5
{ // callseq 11, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd2;
.param .b64 param2;
st.param.b64 [param2+0], %rd3;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf21add_mod_384_uncheckedEPyPKyS2_,
(
param0,
param1,
param2
);
} // callseq 11
.loc 3 846 1
ret;
$L__tmp61:
$L__func_end11:
}
.func _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf11_rshift_384EPyPKy(
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf11_rshift_384EPyPKy_param_0,
.param .b64 _ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf11_rshift_384EPyPKy_param_1
)
{
.reg .b64 %rd<30>;
.loc 3 848 0
$L__func_begin12:
.loc 3 848 0
ld.param.u64 %rd1, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf11_rshift_384EPyPKy_param_0];
ld.param.u64 %rd2, [_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf11_rshift_384EPyPKy_param_1];
$L__tmp62:
.loc 3 849 5
ld.u64 %rd3, [%rd2+8];
shl.b64 %rd4, %rd3, 63;
ld.u64 %rd5, [%rd2];
shr.u64 %rd6, %rd5, 1;
or.b64 %rd7, %rd4, %rd6;
st.u64 [%rd1], %rd7;
.loc 3 850 5
ld.u64 %rd8, [%rd2+16];
shl.b64 %rd9, %rd8, 63;
ld.u64 %rd10, [%rd2+8];
shr.u64 %rd11, %rd10, 1;
or.b64 %rd12, %rd9, %rd11;
st.u64 [%rd1+8], %rd12;
.loc 3 851 5
ld.u64 %rd13, [%rd2+24];
shl.b64 %rd14, %rd13, 63;
ld.u64 %rd15, [%rd2+16];
shr.u64 %rd16, %rd15, 1;
or.b64 %rd17, %rd14, %rd16;
st.u64 [%rd1+16], %rd17;
.loc 3 852 5
ld.u64 %rd18, [%rd2+32];
shl.b64 %rd19, %rd18, 63;
ld.u64 %rd20, [%rd2+24];
shr.u64 %rd21, %rd20, 1;
or.b64 %rd22, %rd19, %rd21;
st.u64 [%rd1+24], %rd22;
.loc 3 853 5
ld.u64 %rd23, [%rd2+40];
shl.b64 %rd24, %rd23, 63;
ld.u64 %rd25, [%rd2+32];
shr.u64 %rd26, %rd25, 1;
or.b64 %rd27, %rd24, %rd26;
st.u64 [%rd1+32], %rd27;
.loc 3 854 5
ld.u64 %rd28, [%rd2+40];
shr.u64 %rd29, %rd28, 1;
st.u64 [%rd1+40], %rd29;
.loc 3 855 1
ret;
$L__tmp63:
$L__func_end12:
}
// .globl _Z16div_by_2_mod_384PyPKy
.visible .func _Z16div_by_2_mod_384PyPKy(
.param .b64 _Z16div_by_2_mod_384PyPKy_param_0,
.param .b64 _Z16div_by_2_mod_384PyPKy_param_1
)
{
.reg .b64 %rd<3>;
.loc 3 857 0
$L__func_begin13:
.loc 3 857 0
ld.param.u64 %rd1, [_Z16div_by_2_mod_384PyPKy_param_0];
ld.param.u64 %rd2, [_Z16div_by_2_mod_384PyPKy_param_1];
$L__tmp64:
.loc 3 858 5
{ // callseq 12, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd1;
.param .b64 param1;
st.param.b64 [param1+0], %rd2;
call.uni
_ZN65_INTERNAL_43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf11_rshift_384EPyPKy,
(
param0,
param1
);
} // callseq 12
.loc 3 859 1
ret;
$L__tmp65:
$L__func_end13:
}
// .globl _Z12cneg_mod_384PyPKybS1_
.visible .func _Z12cneg_mod_384PyPKybS1_(
.param .b64 _Z12cneg_mod_384PyPKybS1__param_0,
.param .b64 _Z12cneg_mod_384PyPKybS1__param_1,
.param .b32 _Z12cneg_mod_384PyPKybS1__param_2,
.param .b64 _Z12cneg_mod_384PyPKybS1__param_3
)
{
.reg .pred %p<5>;
.reg .b16 %rs<3>;
.reg .b64 %rd<12>;
.loc 3 861 0
$L__func_begin14:
.loc 3 861 0
ld.param.u64 %rd5, [_Z12cneg_mod_384PyPKybS1__param_0];
ld.param.s8 %rs1, [_Z12cneg_mod_384PyPKybS1__param_2];
ld.param.u64 %rd6, [_Z12cneg_mod_384PyPKybS1__param_1];
ld.param.u64 %rd7, [_Z12cneg_mod_384PyPKybS1__param_3];
$L__tmp66:
.loc 3 863 5
setp.ne.s16 %p1, %rs1, 0;
not.pred %p2, %p1;
@%p2 bra $L__BB14_2;
bra.uni $L__BB14_1;
$L__BB14_1:
$L__tmp67:
.loc 3 864 9
{ // callseq 13, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd5;
.param .b64 param1;
st.param.b64 [param1+0], %rd7;
.param .b64 param2;
st.param.b64 [param2+0], %rd6;
.param .b64 param3;
st.param.b64 [param3+0], %rd7;
call.uni
_Z11sub_mod_384PyPKyS1_S1_,
(
param0,
param1,
param2,
param3
);
} // callseq 13
bra.uni $L__BB14_5;
$L__tmp68:
$L__BB14_2:
.loc 3 0 9
mov.u64 %rd8, 0;
$L__tmp69:
.loc 3 866 9
mov.pred %p3, 0;
mov.u64 %rd11, %rd8;
@%p3 bra $L__BB14_4;
bra.uni $L__BB14_3;
$L__BB14_3:
.loc 3 0 9
mov.u64 %rd3, %rd11;
add.s64 %rd9, %rd6, %rd3;
ld.u8 %rs2, [%rd9];
add.s64 %rd10, %rd5, %rd3;
st.u8 [%rd10], %rs2;
add.s64 %rd4, %rd3, 1;
setp.lt.u64 %p4, %rd4, 48;
mov.u64 %rd11, %rd4;
@%p4 bra $L__BB14_3;
bra.uni $L__BB14_4;
$L__BB14_4:
.loc 3 866 9
bra.uni $L__BB14_5;
$L__tmp70:
$L__BB14_5:
.loc 3 868 1
ret;
$L__tmp71:
$L__func_end14:
}
.file 1 "/home/p/aleo/snarkVM/algorithms/src/msm/variable_base/blst_377_cuda/./types.h"
.file 2 "/home/p/aleo/snarkVM/algorithms/src/msm/variable_base/blst_377_cuda/./asm_cuda.h"
.file 3 "/home/p/aleo/snarkVM/algorithms/src/msm/variable_base/blst_377_cuda/./asm_cuda.cu"
.section .debug_loc
{
.b64 $L__tmp1
.b64 $L__tmp2
.b8 5
.b8 0
.b8 144
.b8 177
.b8 228
.b8 149
.b8 1
.b64 $L__tmp2
.b64 $L__tmp3
.b8 6
.b8 0
.b8 144
.b8 178
.b8 226
.b8 200
.b8 171
.b8 2
.b64 $L__tmp3
.b64 $L__tmp9
.b8 5
.b8 0
.b8 144
.b8 178
.b8 228
.b8 149
.b8 1
.b64 $L__tmp9
.b64 $L__tmp10
.b8 5
.b8 0
.b8 144
.b8 181
.b8 228
.b8 149
.b8 1
.b64 $L__tmp10
.b64 $L__func_end0
.b8 6
.b8 0
.b8 144
.b8 178
.b8 226
.b8 200
.b8 171
.b8 2
.b64 0
.b64 0
.b64 $L__tmp14
.b64 $L__tmp15
.b8 5
.b8 0
.b8 144
.b8 177
.b8 228
.b8 149
.b8 1
.b64 $L__tmp15
.b64 $L__tmp16
.b8 6
.b8 0
.b8 144
.b8 178
.b8 226
.b8 200
.b8 171
.b8 2
.b64 $L__tmp16
.b64 $L__tmp22
.b8 5
.b8 0
.b8 144
.b8 178
.b8 228
.b8 149
.b8 1
.b64 $L__tmp22
.b64 $L__tmp23
.b8 5
.b8 0
.b8 144
.b8 181
.b8 228
.b8 149
.b8 1
.b64 $L__tmp23
.b64 $L__func_end1
.b8 6
.b8 0
.b8 144
.b8 178
.b8 226
.b8 200
.b8 171
.b8 2
.b64 0
.b64 0
.b64 $L__tmp33
.b64 $L__tmp36
.b8 7
.b8 0
.b8 144
.b8 179
.b8 228
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp36
.b64 $L__tmp38
.b8 7
.b8 0
.b8 144
.b8 181
.b8 232
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp38
.b64 $L__tmp40
.b8 7
.b8 0
.b8 144
.b8 184
.b8 236
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp40
.b64 $L__tmp42
.b8 7
.b8 0
.b8 144
.b8 177
.b8 242
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp42
.b64 $L__tmp44
.b8 8
.b8 0
.b8 144
.b8 180
.b8 226
.b8 196
.b8 161
.b8 166
.b8 174
.b8 9
.b64 $L__tmp44
.b64 $L__func_end4
.b8 8
.b8 0
.b8 144
.b8 183
.b8 230
.b8 196
.b8 161
.b8 166
.b8 174
.b8 9
.b64 0
.b64 0
.b64 $L__tmp34
.b64 $L__tmp35
.b8 8
.b8 0
.b8 144
.b8 183
.b8 232
.b8 196
.b8 161
.b8 166
.b8 174
.b8 9
.b64 $L__tmp35
.b64 $L__tmp37
.b8 6
.b8 0
.b8 144
.b8 185
.b8 200
.b8 201
.b8 171
.b8 2
.b64 $L__tmp37
.b64 $L__tmp39
.b8 7
.b8 0
.b8 144
.b8 177
.b8 230
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp39
.b64 $L__tmp41
.b8 7
.b8 0
.b8 144
.b8 180
.b8 234
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp41
.b64 $L__tmp43
.b8 7
.b8 0
.b8 144
.b8 183
.b8 238
.b8 144
.b8 147
.b8 215
.b8 4
.b64 $L__tmp43
.b64 $L__func_end4
.b8 8
.b8 0
.b8 144
.b8 176
.b8 224
.b8 196
.b8 161
.b8 166
.b8 174
.b8 9
.b64 0
.b64 0
}
.section .debug_abbrev
{
.b8 1
.b8 17
.b8 1
.b8 37
.b8 8
.b8 19
.b8 5
.b8 3
.b8 8
.b8 16
.b8 6
.b8 27
.b8 8
.b8 17
.b8 1
.b8 0
.b8 0
.b8 2
.b8 52
.b8 0
.b8 3
.b8 8
.b8 73
.b8 19
.b8 58
.b8 11
.b8 59
.b8 11
.b8 51
.b8 11
.b8 2
.b8 10
.b8 135,64
.b8 8
.b8 0
.b8 0
.b8 3
.b8 1
.b8 1
.b8 73
.b8 19
.b8 0
.b8 0
.b8 4
.b8 33
.b8 0
.b8 73
.b8 19
.b8 55
.b8 11
.b8 0
.b8 0
.b8 5
.b8 38
.b8 0
.b8 73
.b8 19
.b8 0
.b8 0
.b8 6
.b8 22
.b8 0
.b8 73
.b8 19
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 0
.b8 0
.b8 7
.b8 36
.b8 0
.b8 3
.b8 8
.b8 62
.b8 11
.b8 11
.b8 11
.b8 0
.b8 0
.b8 8
.b8 36
.b8 0
.b8 3
.b8 8
.b8 11
.b8 11
.b8 62
.b8 11
.b8 0
.b8 0
.b8 9
.b8 46
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 135,64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 0
.b8 0
.b8 10
.b8 5
.b8 0
.b8 2
.b8 10
.b8 51
.b8 11
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 0
.b8 0
.b8 11
.b8 11
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 0
.b8 0
.b8 12
.b8 52
.b8 0
.b8 2
.b8 6
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 0
.b8 0
.b8 13
.b8 52
.b8 0
.b8 51
.b8 11
.b8 2
.b8 10
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 73
.b8 19
.b8 0
.b8 0
.b8 14
.b8 46
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 135,64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 5
.b8 73
.b8 19
.b8 63
.b8 12
.b8 0
.b8 0
.b8 15
.b8 5
.b8 0
.b8 2
.b8 10
.b8 51
.b8 11
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 5
.b8 73
.b8 19
.b8 0
.b8 0
.b8 16
.b8 52
.b8 0
.b8 51
.b8 11
.b8 2
.b8 10
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 5
.b8 73
.b8 19
.b8 0
.b8 0
.b8 17
.b8 46
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 135,64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 5
.b8 73
.b8 19
.b8 0
.b8 0
.b8 18
.b8 59
.b8 0
.b8 3
.b8 8
.b8 0
.b8 0
.b8 19
.b8 15
.b8 0
.b8 73
.b8 19
.b8 51
.b8 6
.b8 0
.b8 0
.b8 0
}
.section .debug_info
{
.b32 4062
.b8 2
.b8 0
.b32 .debug_abbrev
.b8 8
.b8 1
.b8 108,103,101,110,102,101,58,32,69,68,71,32,54,46,49
.b8 0
.b8 4
.b8 0
.b8 46,47,97,115,109,95,99,117,100,97,46,99,117
.b8 0
.b32 .debug_line
.b8 47,104,111,109,101,47,112,47,97,108,101,111,47,115,110,97,114,107,86,77,47,97,108,103,111,114,105,116,104,109,115,47,115,114,99,47,109,115,109,47
.b8 118,97,114,105,97,98,108,101,95,98,97,115,101,47,98,108,115,116,95,51,55,55,95,99,117,100,97
.b8 0
.b64 0
.b8 2
.b8 66,76,83,49,50,95,51,55,55,95,80
.b8 0
.b32 238
.b8 1
.b8 46
.b8 5
.b8 9
.b8 3
.b64 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_P
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,54,95,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48
.b8 48,48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,95,66,76,83,49,50,95,51,55
.b8 55,95,80
.b8 0
.b8 3
.b32 250
.b8 4
.b32 291
.b8 6
.b8 0
.b8 5
.b32 255
.b8 6
.b32 269
.b8 108,105,109,98,95,116
.b8 0
.b8 1
.b8 11
.b8 7
.b8 117,110,115,105,103,110,101,100,32,108,111,110,103,32,108,111,110,103
.b8 0
.b8 7
.b8 8
.b8 8
.b8 95,95,65,82,82,65,89,95,83,73,90,69,95,84,89,80,69,95,95
.b8 0
.b8 8
.b8 7
.b8 2
.b8 66,76,83,49,50,95,51,55,55,95,79,78,69
.b8 0
.b32 238
.b8 1
.b8 51
.b8 5
.b8 9
.b8 3
.b64 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_ONE
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,54,95,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48
.b8 48,48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,95,66,76,83,49,50,95,51,55
.b8 55,95,79,78,69
.b8 0
.b8 2
.b8 66,76,83,49,50,95,51,55,55,95,82,50
.b8 0
.b32 238
.b8 1
.b8 52
.b8 5
.b8 9
.b8 3
.b64 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_R2
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,54,95,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48
.b8 48,48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,95,66,76,83,49,50,95,51,55
.b8 55,95,82,50
.b8 0
.b8 2
.b8 66,76,83,49,50,95,51,55,55,95,112,48
.b8 0
.b32 250
.b8 1
.b8 60
.b8 5
.b8 9
.b8 3
.b64 __nv_static_56__43_tmpxft_0028c5ce_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_p0
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,54,95,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48
.b8 48,48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,95,66,76,83,49,50,95,51,55
.b8 55,95,112,48
.b8 0
.b8 6
.b32 679
.b8 98,108,115,116,95,102,112
.b8 0
.b8 1
.b8 32
.b8 3
.b32 255
.b8 4
.b32 291
.b8 6
.b8 0
.b8 9
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,57,105,115,95,103,116,95,51,56,52
.b8 69,80,75,121,83,49,95
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,57,105,115,95,103,116,95,51,56,52
.b8 69,80,75,121,83,49,95
.b8 0
.b8 2
.b8 21
.b32 4014
.b8 10
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 108,101,102,116
.b8 0
.b8 2
.b8 21
.b32 4027
.b8 10
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,105,103,104,116
.b8 0
.b8 2
.b8 21
.b32 4027
.b8 11
.b64 $L__tmp0
.b64 $L__tmp11
.b8 12
.b32 .debug_loc
.b8 105
.b8 0
.b8 2
.b8 22
.b32 4014
.b8 0
.b8 0
.b8 9
.b64 $L__func_begin1
.b64 $L__func_end1
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,57,105,115,95,103,101,95,51,56,52
.b8 69,80,75,121,83,49,95
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,57,105,115,95,103,101,95,51,56,52
.b8 69,80,75,121,83,49,95
.b8 0
.b8 3
.b8 7
.b32 4014
.b8 10
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 108,101,102,116
.b8 0
.b8 3
.b8 7
.b32 4027
.b8 10
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,105,103,104,116
.b8 0
.b8 3
.b8 7
.b32 4027
.b8 11
.b64 $L__tmp13
.b64 $L__tmp24
.b8 12
.b32 .debug_loc+133
.b8 105
.b8 0
.b8 3
.b8 8
.b32 4014
.b8 0
.b8 0
.b8 9
.b64 $L__func_begin2
.b64 $L__func_end2
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,50,49,115,117,98,95,109,111,100,95
.b8 51,56,52,95,117,110,99,104,101,99,107,101,100,69,80,121,80,75,121,83,50,95
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,50,49,115,117,98,95,109,111,100,95
.b8 51,56,52,95,117,110,99,104,101,99,107,101,100,69,80,121,80,75,121,83,50,95
.b8 0
.b8 3
.b8 18
.b32 4021
.b8 10
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 18
.b32 4036
.b8 10
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 18
.b32 4027
.b8 10
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 18
.b32 4027
.b8 0
.b8 9
.b64 $L__func_begin3
.b64 $L__func_end3
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,54,114,101,100,117,99,101,69,80,121
.b8 80,75,121
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,54,114,101,100,117,99,101,69,80,121
.b8 80,75,121
.b8 0
.b8 3
.b8 48
.b32 4021
.b8 10
.b8 6
.b8 144
.b8 181
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 120
.b8 0
.b8 3
.b8 48
.b32 4036
.b8 10
.b8 6
.b8 144
.b8 182
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 48
.b32 4027
.b8 11
.b64 $L__tmp29
.b64 $L__tmp30
.b8 13
.b8 6
.b8 11
.b8 3
.b64 __local_depot3
.b8 35
.b8 0
.b8 120,95,115,117,98
.b8 0
.b8 3
.b8 50
.b32 664
.b8 0
.b8 0
.b8 9
.b64 $L__func_begin4
.b64 $L__func_end4
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,56,109,111,110,116,95,51,56,52,69
.b8 80,121,83,48,95,80,75,121,121
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,56,109,111,110,116,95,51,56,52,69
.b8 80,121,83,48,95,80,75,121,121
.b8 0
.b8 3
.b8 59
.b32 4021
.b8 10
.b8 6
.b8 144
.b8 181
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 59
.b32 4036
.b8 10
.b8 6
.b8 144
.b8 182
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114
.b8 0
.b8 3
.b8 59
.b32 4036
.b8 10
.b8 6
.b8 144
.b8 183
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 59
.b32 4027
.b8 10
.b8 6
.b8 144
.b8 184
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112,95,105,110,118
.b8 0
.b8 3
.b8 59
.b32 250
.b8 11
.b64 $L__tmp32
.b64 $L__tmp45
.b8 12
.b32 .debug_loc+266
.b8 107
.b8 0
.b8 3
.b8 61
.b32 255
.b8 12
.b32 .debug_loc+434
.b8 99,114,111,115,115,95,99,97,114,114,121
.b8 0
.b8 3
.b8 63
.b32 255
.b8 0
.b8 0
.b8 14
.b64 $L__func_begin5
.b64 $L__func_end5
.b8 1
.b8 156
.b8 95,90,49,50,109,117,108,95,109,111,110,116,95,51,56,52,80,121,80,75,121,83,49,95,83,49,95,121
.b8 0
.b8 95,90,49,50,109,117,108,95,109,111,110,116,95,51,56,52,80,121,80,75,121,83,49,95,83,49,95,121
.b8 0
.b8 3
.b8 166
.b8 1
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 166
.b8 1
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 166
.b8 1
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 166
.b8 1
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 180
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 166
.b8 1
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 181
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112,95,105,110,118
.b8 0
.b8 3
.b8 166
.b8 1
.b32 255
.b8 11
.b64 $L__tmp46
.b64 $L__tmp47
.b8 16
.b8 6
.b8 11
.b8 3
.b64 __local_depot5
.b8 35
.b8 0
.b8 114
.b8 0
.b8 3
.b8 167
.b8 1
.b32 4045
.b8 0
.b8 0
.b8 14
.b64 $L__func_begin6
.b64 $L__func_end6
.b8 1
.b8 156
.b8 95,90,49,50,115,113,114,95,109,111,110,116,95,51,56,52,80,121,80,75,121,83,49,95,121
.b8 0
.b8 95,90,49,50,115,113,114,95,109,111,110,116,95,51,56,52,80,121,80,75,121,83,49,95,121
.b8 0
.b8 3
.b8 116
.b8 2
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 116
.b8 2
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 116
.b8 2
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 116
.b8 2
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 180
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112,95,105,110,118
.b8 0
.b8 3
.b8 116
.b8 2
.b32 255
.b8 11
.b64 $L__tmp48
.b64 $L__tmp49
.b8 16
.b8 6
.b8 11
.b8 3
.b64 __local_depot6
.b8 35
.b8 0
.b8 114
.b8 0
.b8 3
.b8 117
.b8 2
.b32 4045
.b8 0
.b8 0
.b8 17
.b64 $L__func_begin7
.b64 $L__func_end7
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,50,49,97,100,100,95,109,111,100,95
.b8 51,56,52,95,117,110,99,104,101,99,107,101,100,69,80,121,80,75,121,83,50,95
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,50,49,97,100,100,95,109,111,100,95
.b8 51,56,52,95,117,110,99,104,101,99,107,101,100,69,80,121,80,75,121,83,50,95
.b8 0
.b8 3
.b8 17
.b8 3
.b32 4021
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 17
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 17
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 17
.b8 3
.b32 4027
.b8 0
.b8 14
.b64 $L__func_begin8
.b64 $L__func_end8
.b8 1
.b8 156
.b8 95,90,49,49,97,100,100,95,109,111,100,95,51,56,52,80,121,80,75,121,83,49,95,83,49,95
.b8 0
.b8 95,90,49,49,97,100,100,95,109,111,100,95,51,56,52,80,121,80,75,121,83,49,95,83,49,95
.b8 0
.b8 3
.b8 47
.b8 3
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 47
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 47
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 47
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 180
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 47
.b8 3
.b32 4027
.b8 0
.b8 14
.b64 $L__func_begin9
.b64 $L__func_end9
.b8 1
.b8 156
.b8 95,90,49,49,115,117,98,95,109,111,100,95,51,56,52,80,121,80,75,121,83,49,95,83,49,95
.b8 0
.b8 95,90,49,49,115,117,98,95,109,111,100,95,51,56,52,80,121,80,75,121,83,49,95,83,49,95
.b8 0
.b8 3
.b8 54
.b8 3
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 181
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 54
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 182
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 54
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 183
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 54
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 184
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 54
.b8 3
.b32 4027
.b8 11
.b64 $L__tmp54
.b64 $L__tmp57
.b8 16
.b8 6
.b8 11
.b8 3
.b64 __local_depot9
.b8 35
.b8 0
.b8 97,100,100,101,100
.b8 0
.b8 3
.b8 55
.b8 3
.b32 664
.b8 0
.b8 0
.b8 14
.b64 $L__func_begin10
.b64 $L__func_end10
.b8 1
.b8 156
.b8 95,90,49,56,115,117,98,95,109,111,100,95,51,56,52,95,117,110,115,97,102,101,80,121,80,75,121,83,49,95
.b8 0
.b8 95,90,49,56,115,117,98,95,109,111,100,95,51,56,52,95,117,110,115,97,102,101,80,121,80,75,121,83,49,95
.b8 0
.b8 3
.b8 70
.b8 3
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 70
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 70
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 70
.b8 3
.b32 4027
.b8 0
.b8 14
.b64 $L__func_begin11
.b64 $L__func_end11
.b8 1
.b8 156
.b8 95,90,49,56,97,100,100,95,109,111,100,95,51,56,52,95,117,110,115,97,102,101,80,121,80,75,121,83,49,95
.b8 0
.b8 95,90,49,56,97,100,100,95,109,111,100,95,51,56,52,95,117,110,115,97,102,101,80,121,80,75,121,83,49,95
.b8 0
.b8 3
.b8 75
.b8 3
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 75
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 75
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 179
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 98
.b8 0
.b8 3
.b8 75
.b8 3
.b32 4027
.b8 0
.b8 17
.b64 $L__func_begin12
.b64 $L__func_end12
.b8 1
.b8 156
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,49,49,95,114,115,104,105,102,116,95
.b8 51,56,52,69,80,121,80,75,121
.b8 0
.b8 95,90,78,54,53,95,73,78,84,69,82,78,65,76,95,52,51,95,116,109,112,120,102,116,95,48,48,50,56,99,53,99,101,95,48,48,48,48,48,48
.b8 48,48,95,55,95,97,115,109,95,99,117,100,97,95,99,112,112,49,95,105,105,95,55,49,53,97,100,56,99,102,49,49,95,114,115,104,105,102,116,95
.b8 51,56,52,69,80,121,80,75,121
.b8 0
.b8 3
.b8 80
.b8 3
.b32 4021
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 80
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 118,97,108,117,101
.b8 0
.b8 3
.b8 80
.b8 3
.b32 4027
.b8 0
.b8 14
.b64 $L__func_begin13
.b64 $L__func_end13
.b8 1
.b8 156
.b8 95,90,49,54,100,105,118,95,98,121,95,50,95,109,111,100,95,51,56,52,80,121,80,75,121
.b8 0
.b8 95,90,49,54,100,105,118,95,98,121,95,50,95,109,111,100,95,51,56,52,80,121,80,75,121
.b8 0
.b8 3
.b8 89
.b8 3
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 89
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 89
.b8 3
.b32 4027
.b8 0
.b8 14
.b64 $L__func_begin14
.b64 $L__func_end14
.b8 1
.b8 156
.b8 95,90,49,50,99,110,101,103,95,109,111,100,95,51,56,52,80,121,80,75,121,98,83,49,95
.b8 0
.b8 95,90,49,50,99,110,101,103,95,109,111,100,95,51,56,52,80,121,80,75,121,98,83,49,95
.b8 0
.b8 3
.b8 93
.b8 3
.b32 4021
.b8 1
.b8 15
.b8 6
.b8 144
.b8 181
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 114,101,116
.b8 0
.b8 3
.b8 93
.b8 3
.b32 4036
.b8 15
.b8 6
.b8 144
.b8 182
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 97
.b8 0
.b8 3
.b8 93
.b8 3
.b32 4027
.b8 15
.b8 6
.b8 144
.b8 177
.b8 230
.b8 201
.b8 171
.b8 2
.b8 2
.b8 102,108,97,103
.b8 0
.b8 3
.b8 93
.b8 3
.b32 4057
.b8 15
.b8 6
.b8 144
.b8 183
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 112
.b8 0
.b8 3
.b8 93
.b8 3
.b32 4027
.b8 0
.b8 7
.b8 105,110,116
.b8 0
.b8 5
.b8 4
.b8 18
.b8 118,111,105,100
.b8 0
.b8 19
.b32 250
.b32 12
.b8 19
.b32 255
.b32 12
.b8 3
.b32 255
.b8 4
.b32 291
.b8 12
.b8 0
.b8 7
.b8 98,111,111,108
.b8 0
.b8 2
.b8 1
.b8 0
}
.section .debug_macinfo
{
.b8 0
}