//
// 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
.address_size 64
// .globl _Z12mul_mont_384PyPKyS1_S1_y
.global .align 8 .b8 __nv_static_56__43_tmpxft_0028c598_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_0028c598_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_0028c598_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_0028c598_00000000_7_asm_cuda_cpp1_ii_715ad8cf_BLS12_377_p0 = -8860621160618917889;
.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
)
{
.reg .pred %p<12>;
.reg .b64 %rd<305>;
ld.param.u64 %rd12, [_Z12mul_mont_384PyPKyS1_S1_y_param_0];
ld.param.u64 %rd186, [_Z12mul_mont_384PyPKyS1_S1_y_param_1];
ld.param.u64 %rd187, [_Z12mul_mont_384PyPKyS1_S1_y_param_2];
ld.param.u64 %rd13, [_Z12mul_mont_384PyPKyS1_S1_y_param_3];
ld.param.u64 %rd188, [_Z12mul_mont_384PyPKyS1_S1_y_param_4];
ld.u64 %rd26, [%rd186];
ld.u64 %rd27, [%rd186+8];
ld.u64 %rd28, [%rd186+16];
ld.u64 %rd29, [%rd186+24];
ld.u64 %rd30, [%rd186+32];
ld.u64 %rd31, [%rd186+40];
ld.u64 %rd32, [%rd187];
ld.u64 %rd33, [%rd187+8];
ld.u64 %rd34, [%rd187+16];
ld.u64 %rd35, [%rd187+24];
ld.u64 %rd36, [%rd187+32];
ld.u64 %rd37, [%rd187+40];
// begin inline asm
{
.reg .u64 c;
.reg .u64 nc;
.reg .u64 t;
mad.lo.cc.u64 %rd50, %rd26, %rd32, 0;
madc.hi.cc.u64 c, %rd26, %rd32, 0;
madc.lo.cc.u64 %rd72, %rd26, %rd33, c;
madc.hi.cc.u64 c, %rd26, %rd33, 0;
madc.lo.cc.u64 %rd73, %rd26, %rd34, c;
madc.hi.cc.u64 c, %rd26, %rd34, 0;
madc.lo.cc.u64 %rd74, %rd26, %rd35, c;
madc.hi.cc.u64 c, %rd26, %rd35, 0;
madc.lo.cc.u64 %rd75, %rd26, %rd36, c;
madc.hi.cc.u64 c, %rd26, %rd36, 0;
madc.lo.cc.u64 %rd76, %rd26, %rd37, c;
madc.hi.u64 %rd77, %rd26, %rd37, 0;
mad.lo.cc.u64 %rd72, %rd27, %rd32, %rd72;
madc.hi.cc.u64 c, %rd27, %rd32, 0;
addc.cc.u64 t, %rd73, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd73, %rd27, %rd33, t;
madc.hi.cc.u64 c, %rd27, %rd33, nc;
addc.cc.u64 t, %rd74, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd74, %rd27, %rd34, t;
madc.hi.cc.u64 c, %rd27, %rd34, nc;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd27, %rd35, t;
madc.hi.cc.u64 c, %rd27, %rd35, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd27, %rd36, t;
madc.hi.cc.u64 c, %rd27, %rd36, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd27, %rd37, t;
madc.hi.u64 %rd100, %rd27, %rd37, nc;
mad.lo.cc.u64 %rd73, %rd28, %rd32, %rd73;
madc.hi.cc.u64 c, %rd28, %rd32, 0;
addc.cc.u64 t, %rd74, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd74, %rd28, %rd33, t;
madc.hi.cc.u64 c, %rd28, %rd33, nc;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd28, %rd34, t;
madc.hi.cc.u64 c, %rd28, %rd34, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd28, %rd35, t;
madc.hi.cc.u64 c, %rd28, %rd35, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd28, %rd36, t;
madc.hi.cc.u64 c, %rd28, %rd36, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd28, %rd37, t;
madc.hi.u64 %rd123, %rd28, %rd37, nc;
mad.lo.cc.u64 %rd74, %rd29, %rd32, %rd74;
madc.hi.cc.u64 c, %rd29, %rd32, 0;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd29, %rd33, t;
madc.hi.cc.u64 c, %rd29, %rd33, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd29, %rd34, t;
madc.hi.cc.u64 c, %rd29, %rd34, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd29, %rd35, t;
madc.hi.cc.u64 c, %rd29, %rd35, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd29, %rd36, t;
madc.hi.cc.u64 c, %rd29, %rd36, nc;
addc.cc.u64 t, %rd123, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd123, %rd29, %rd37, t;
madc.hi.u64 %rd146, %rd29, %rd37, nc;
mad.lo.cc.u64 %rd75, %rd30, %rd32, %rd75;
madc.hi.cc.u64 c, %rd30, %rd32, 0;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd30, %rd33, t;
madc.hi.cc.u64 c, %rd30, %rd33, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd30, %rd34, t;
madc.hi.cc.u64 c, %rd30, %rd34, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd30, %rd35, t;
madc.hi.cc.u64 c, %rd30, %rd35, nc;
addc.cc.u64 t, %rd123, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd123, %rd30, %rd36, t;
madc.hi.cc.u64 c, %rd30, %rd36, nc;
addc.cc.u64 t, %rd146, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd146, %rd30, %rd37, t;
madc.hi.u64 %rd169, %rd30, %rd37, nc;
mad.lo.cc.u64 %rd76, %rd31, %rd32, %rd76;
madc.hi.cc.u64 c, %rd31, %rd32, 0;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd31, %rd33, t;
madc.hi.cc.u64 c, %rd31, %rd33, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd31, %rd34, t;
madc.hi.cc.u64 c, %rd31, %rd34, nc;
addc.cc.u64 t, %rd123, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd123, %rd31, %rd35, t;
madc.hi.cc.u64 c, %rd31, %rd35, nc;
addc.cc.u64 t, %rd146, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd146, %rd31, %rd36, t;
madc.hi.cc.u64 c, %rd31, %rd36, nc;
addc.cc.u64 t, %rd169, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd169, %rd31, %rd37, t;
madc.hi.u64 %rd170, %rd31, %rd37, nc;
}
// end inline asm
mul.lo.s64 %rd64, %rd50, %rd188;
ld.u64 %rd58, [%rd13];
ld.u64 %rd59, [%rd13+8];
ld.u64 %rd60, [%rd13+16];
ld.u64 %rd61, [%rd13+24];
ld.u64 %rd62, [%rd13+32];
ld.u64 %rd63, [%rd13+40];
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd64, %rd58, %rd50;
madc.hi.cc.u64 c, %rd64, %rd58, 0;
addc.cc.u64 t, %rd72, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd72, %rd64, %rd59, t;
madc.hi.cc.u64 c, %rd64, %rd59, nc;
addc.cc.u64 t, %rd73, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd73, %rd64, %rd60, t;
madc.hi.cc.u64 c, %rd64, %rd60, nc;
addc.cc.u64 t, %rd74, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd74, %rd64, %rd61, t;
madc.hi.cc.u64 c, %rd64, %rd61, nc;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd64, %rd62, t;
madc.hi.cc.u64 c, %rd64, %rd62, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd64, %rd63, t;
madc.hi.cc.u64 c, %rd64, %rd63, nc;
addc.cc.u64 %rd77, %rd77, c;
addc.u64 %rd102, 0, 0;
}
// end inline asm
mul.lo.s64 %rd86, %rd72, %rd188;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd86, %rd58, %rd72;
madc.hi.cc.u64 c, %rd86, %rd58, 0;
addc.cc.u64 t, %rd73, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd73, %rd86, %rd59, t;
madc.hi.cc.u64 c, %rd86, %rd59, nc;
addc.cc.u64 t, %rd74, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd74, %rd86, %rd60, t;
madc.hi.cc.u64 c, %rd86, %rd60, nc;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd86, %rd61, t;
madc.hi.cc.u64 c, %rd86, %rd61, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd86, %rd62, t;
madc.hi.cc.u64 c, %rd86, %rd62, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd86, %rd63, t;
madc.hi.cc.u64 c, %rd86, %rd63, nc;
addc.cc.u64 c, c, %rd102;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd100, %rd100, c;
addc.u64 %rd102, nc, 0;
}
// end inline asm
mul.lo.s64 %rd109, %rd73, %rd188;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd109, %rd58, %rd73;
madc.hi.cc.u64 c, %rd109, %rd58, 0;
addc.cc.u64 t, %rd74, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd74, %rd109, %rd59, t;
madc.hi.cc.u64 c, %rd109, %rd59, nc;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd109, %rd60, t;
madc.hi.cc.u64 c, %rd109, %rd60, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd109, %rd61, t;
madc.hi.cc.u64 c, %rd109, %rd61, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd109, %rd62, t;
madc.hi.cc.u64 c, %rd109, %rd62, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd109, %rd63, t;
madc.hi.cc.u64 c, %rd109, %rd63, nc;
addc.cc.u64 c, c, %rd102;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd123, %rd123, c;
addc.u64 %rd102, nc, 0;
}
// end inline asm
mul.lo.s64 %rd132, %rd74, %rd188;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd132, %rd58, %rd74;
madc.hi.cc.u64 c, %rd132, %rd58, 0;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd132, %rd59, t;
madc.hi.cc.u64 c, %rd132, %rd59, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd132, %rd60, t;
madc.hi.cc.u64 c, %rd132, %rd60, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd132, %rd61, t;
madc.hi.cc.u64 c, %rd132, %rd61, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd132, %rd62, t;
madc.hi.cc.u64 c, %rd132, %rd62, nc;
addc.cc.u64 t, %rd123, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd123, %rd132, %rd63, t;
madc.hi.cc.u64 c, %rd132, %rd63, nc;
addc.cc.u64 c, c, %rd102;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd146, %rd146, c;
addc.u64 %rd102, nc, 0;
}
// end inline asm
mul.lo.s64 %rd155, %rd75, %rd188;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd155, %rd58, %rd75;
madc.hi.cc.u64 c, %rd155, %rd58, 0;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd155, %rd59, t;
madc.hi.cc.u64 c, %rd155, %rd59, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd155, %rd60, t;
madc.hi.cc.u64 c, %rd155, %rd60, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd155, %rd61, t;
madc.hi.cc.u64 c, %rd155, %rd61, nc;
addc.cc.u64 t, %rd123, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd123, %rd155, %rd62, t;
madc.hi.cc.u64 c, %rd155, %rd62, nc;
addc.cc.u64 t, %rd146, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd146, %rd155, %rd63, t;
madc.hi.cc.u64 c, %rd155, %rd63, nc;
addc.cc.u64 c, c, %rd102;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd169, %rd169, c;
addc.u64 %rd102, nc, 0;
}
// end inline asm
mul.lo.s64 %rd178, %rd76, %rd188;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd178, %rd58, %rd76;
madc.hi.cc.u64 c, %rd178, %rd58, 0;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd178, %rd59, t;
madc.hi.cc.u64 c, %rd178, %rd59, nc;
addc.cc.u64 t, %rd100, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd100, %rd178, %rd60, t;
madc.hi.cc.u64 c, %rd178, %rd60, nc;
addc.cc.u64 t, %rd123, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd123, %rd178, %rd61, t;
madc.hi.cc.u64 c, %rd178, %rd61, nc;
addc.cc.u64 t, %rd146, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd146, %rd178, %rd62, t;
madc.hi.cc.u64 c, %rd178, %rd62, nc;
addc.cc.u64 t, %rd169, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd169, %rd178, %rd63, t;
madc.hi.cc.u64 c, %rd178, %rd63, nc;
addc.cc.u64 c, c, %rd102;
add.u64 %rd170, %rd170, c;
}
// end inline asm
st.u8 [%rd12], %rd77;
shr.u64 %rd201, %rd77, 56;
st.u8 [%rd12+7], %rd201;
shr.u64 %rd202, %rd77, 48;
st.u8 [%rd12+6], %rd202;
shr.u64 %rd203, %rd77, 40;
st.u8 [%rd12+5], %rd203;
shr.u64 %rd204, %rd77, 32;
st.u8 [%rd12+4], %rd204;
shr.u64 %rd205, %rd77, 24;
st.u8 [%rd12+3], %rd205;
shr.u64 %rd206, %rd77, 16;
st.u8 [%rd12+2], %rd206;
shr.u64 %rd207, %rd77, 8;
st.u8 [%rd12+1], %rd207;
st.u8 [%rd12+8], %rd100;
shr.u64 %rd208, %rd100, 56;
st.u8 [%rd12+15], %rd208;
shr.u64 %rd209, %rd100, 48;
st.u8 [%rd12+14], %rd209;
shr.u64 %rd210, %rd100, 40;
st.u8 [%rd12+13], %rd210;
shr.u64 %rd211, %rd100, 32;
st.u8 [%rd12+12], %rd211;
shr.u64 %rd212, %rd100, 24;
st.u8 [%rd12+11], %rd212;
shr.u64 %rd213, %rd100, 16;
st.u8 [%rd12+10], %rd213;
shr.u64 %rd214, %rd100, 8;
st.u8 [%rd12+9], %rd214;
st.u8 [%rd12+16], %rd123;
shr.u64 %rd215, %rd123, 56;
st.u8 [%rd12+23], %rd215;
shr.u64 %rd216, %rd123, 48;
st.u8 [%rd12+22], %rd216;
shr.u64 %rd217, %rd123, 40;
st.u8 [%rd12+21], %rd217;
shr.u64 %rd218, %rd123, 32;
st.u8 [%rd12+20], %rd218;
shr.u64 %rd219, %rd123, 24;
st.u8 [%rd12+19], %rd219;
shr.u64 %rd220, %rd123, 16;
st.u8 [%rd12+18], %rd220;
shr.u64 %rd221, %rd123, 8;
st.u8 [%rd12+17], %rd221;
st.u8 [%rd12+24], %rd146;
shr.u64 %rd222, %rd146, 56;
st.u8 [%rd12+31], %rd222;
shr.u64 %rd223, %rd146, 48;
st.u8 [%rd12+30], %rd223;
shr.u64 %rd224, %rd146, 40;
st.u8 [%rd12+29], %rd224;
shr.u64 %rd225, %rd146, 32;
st.u8 [%rd12+28], %rd225;
shr.u64 %rd226, %rd146, 24;
st.u8 [%rd12+27], %rd226;
shr.u64 %rd227, %rd146, 16;
st.u8 [%rd12+26], %rd227;
shr.u64 %rd228, %rd146, 8;
st.u8 [%rd12+25], %rd228;
st.u8 [%rd12+32], %rd169;
shr.u64 %rd229, %rd169, 56;
st.u8 [%rd12+39], %rd229;
shr.u64 %rd230, %rd169, 48;
st.u8 [%rd12+38], %rd230;
shr.u64 %rd231, %rd169, 40;
st.u8 [%rd12+37], %rd231;
shr.u64 %rd232, %rd169, 32;
st.u8 [%rd12+36], %rd232;
shr.u64 %rd233, %rd169, 24;
st.u8 [%rd12+35], %rd233;
shr.u64 %rd234, %rd169, 16;
st.u8 [%rd12+34], %rd234;
shr.u64 %rd235, %rd169, 8;
st.u8 [%rd12+33], %rd235;
st.u8 [%rd12+40], %rd170;
shr.u64 %rd236, %rd170, 56;
st.u8 [%rd12+47], %rd236;
shr.u64 %rd237, %rd170, 48;
st.u8 [%rd12+46], %rd237;
shr.u64 %rd238, %rd170, 40;
st.u8 [%rd12+45], %rd238;
shr.u64 %rd239, %rd170, 32;
st.u8 [%rd12+44], %rd239;
shr.u64 %rd240, %rd170, 24;
st.u8 [%rd12+43], %rd240;
shr.u64 %rd241, %rd170, 16;
st.u8 [%rd12+42], %rd241;
shr.u64 %rd242, %rd170, 8;
st.u8 [%rd12+41], %rd242;
ld.u64 %rd7, [%rd13+40];
setp.lt.u64 %p1, %rd170, %rd7;
@%p1 bra $L__BB0_12;
setp.gt.u64 %p2, %rd170, %rd7;
ld.u64 %rd8, [%rd13+32];
@%p2 bra $L__BB0_11;
setp.lt.u64 %p3, %rd169, %rd8;
@%p3 bra $L__BB0_12;
setp.gt.u64 %p4, %rd169, %rd8;
@%p4 bra $L__BB0_11;
ld.u64 %rd9, [%rd13+24];
setp.lt.u64 %p5, %rd146, %rd9;
@%p5 bra $L__BB0_12;
setp.gt.u64 %p6, %rd146, %rd9;
@%p6 bra $L__BB0_11;
ld.u64 %rd10, [%rd13+16];
setp.lt.u64 %p7, %rd123, %rd10;
@%p7 bra $L__BB0_12;
setp.gt.u64 %p8, %rd123, %rd10;
@%p8 bra $L__BB0_11;
ld.u64 %rd11, [%rd13+8];
setp.lt.u64 %p9, %rd100, %rd11;
@%p9 bra $L__BB0_12;
setp.gt.u64 %p10, %rd100, %rd11;
@%p10 bra $L__BB0_11;
ld.u64 %rd243, [%rd13];
setp.lt.u64 %p11, %rd77, %rd243;
@%p11 bra $L__BB0_12;
$L__BB0_11:
ld.param.u64 %rd304, [_Z12mul_mont_384PyPKyS1_S1_y_param_0];
ld.u64 %rd256, [%rd13];
ld.u64 %rd257, [%rd13+8];
ld.u64 %rd258, [%rd13+16];
ld.u64 %rd259, [%rd13+24];
// begin inline asm
sub.cc.u64 %rd244, %rd77, %rd256;
subc.cc.u64 %rd245, %rd100, %rd257;
subc.cc.u64 %rd246, %rd123, %rd258;
subc.cc.u64 %rd247, %rd146, %rd259;
subc.cc.u64 %rd248, %rd169, %rd8;
subc.u64 %rd249, %rd170, %rd7;
// end inline asm
st.u8 [%rd304], %rd244;
shr.u64 %rd262, %rd244, 56;
st.u8 [%rd304+7], %rd262;
shr.u64 %rd263, %rd244, 48;
st.u8 [%rd304+6], %rd263;
shr.u64 %rd264, %rd244, 40;
st.u8 [%rd304+5], %rd264;
shr.u64 %rd265, %rd244, 32;
st.u8 [%rd304+4], %rd265;
shr.u64 %rd266, %rd244, 24;
st.u8 [%rd304+3], %rd266;
shr.u64 %rd267, %rd244, 16;
st.u8 [%rd304+2], %rd267;
shr.u64 %rd268, %rd244, 8;
st.u8 [%rd304+1], %rd268;
st.u8 [%rd304+8], %rd245;
shr.u64 %rd269, %rd245, 56;
st.u8 [%rd304+15], %rd269;
shr.u64 %rd270, %rd245, 48;
st.u8 [%rd304+14], %rd270;
shr.u64 %rd271, %rd245, 40;
st.u8 [%rd304+13], %rd271;
shr.u64 %rd272, %rd245, 32;
st.u8 [%rd304+12], %rd272;
shr.u64 %rd273, %rd245, 24;
st.u8 [%rd304+11], %rd273;
shr.u64 %rd274, %rd245, 16;
st.u8 [%rd304+10], %rd274;
shr.u64 %rd275, %rd245, 8;
st.u8 [%rd304+9], %rd275;
st.u8 [%rd304+16], %rd246;
shr.u64 %rd276, %rd246, 56;
st.u8 [%rd304+23], %rd276;
shr.u64 %rd277, %rd246, 48;
st.u8 [%rd304+22], %rd277;
shr.u64 %rd278, %rd246, 40;
st.u8 [%rd304+21], %rd278;
shr.u64 %rd279, %rd246, 32;
st.u8 [%rd304+20], %rd279;
shr.u64 %rd280, %rd246, 24;
st.u8 [%rd304+19], %rd280;
shr.u64 %rd281, %rd246, 16;
st.u8 [%rd304+18], %rd281;
shr.u64 %rd282, %rd246, 8;
st.u8 [%rd304+17], %rd282;
st.u8 [%rd304+24], %rd247;
shr.u64 %rd283, %rd247, 56;
st.u8 [%rd304+31], %rd283;
shr.u64 %rd284, %rd247, 48;
st.u8 [%rd304+30], %rd284;
shr.u64 %rd285, %rd247, 40;
st.u8 [%rd304+29], %rd285;
shr.u64 %rd286, %rd247, 32;
st.u8 [%rd304+28], %rd286;
shr.u64 %rd287, %rd247, 24;
st.u8 [%rd304+27], %rd287;
shr.u64 %rd288, %rd247, 16;
st.u8 [%rd304+26], %rd288;
shr.u64 %rd289, %rd247, 8;
st.u8 [%rd304+25], %rd289;
st.u8 [%rd304+32], %rd248;
shr.u64 %rd290, %rd248, 56;
st.u8 [%rd304+39], %rd290;
shr.u64 %rd291, %rd248, 48;
st.u8 [%rd304+38], %rd291;
shr.u64 %rd292, %rd248, 40;
st.u8 [%rd304+37], %rd292;
shr.u64 %rd293, %rd248, 32;
st.u8 [%rd304+36], %rd293;
shr.u64 %rd294, %rd248, 24;
st.u8 [%rd304+35], %rd294;
shr.u64 %rd295, %rd248, 16;
st.u8 [%rd304+34], %rd295;
shr.u64 %rd296, %rd248, 8;
st.u8 [%rd304+33], %rd296;
st.u8 [%rd304+40], %rd249;
shr.u64 %rd297, %rd249, 56;
st.u8 [%rd304+47], %rd297;
shr.u64 %rd298, %rd249, 48;
st.u8 [%rd304+46], %rd298;
shr.u64 %rd299, %rd249, 40;
st.u8 [%rd304+45], %rd299;
shr.u64 %rd300, %rd249, 32;
st.u8 [%rd304+44], %rd300;
shr.u64 %rd301, %rd249, 24;
st.u8 [%rd304+43], %rd301;
shr.u64 %rd302, %rd249, 16;
st.u8 [%rd304+42], %rd302;
shr.u64 %rd303, %rd249, 8;
st.u8 [%rd304+41], %rd303;
$L__BB0_12:
ret;
}
// .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
)
{
.reg .pred %p<12>;
.reg .b64 %rd<346>;
ld.param.u64 %rd12, [_Z12sqr_mont_384PyPKyS1_y_param_0];
ld.param.u64 %rd210, [_Z12sqr_mont_384PyPKyS1_y_param_1];
ld.param.u64 %rd13, [_Z12sqr_mont_384PyPKyS1_y_param_2];
ld.param.u64 %rd211, [_Z12sqr_mont_384PyPKyS1_y_param_3];
ld.u64 %rd26, [%rd210];
ld.u64 %rd27, [%rd210+8];
ld.u64 %rd28, [%rd210+16];
ld.u64 %rd29, [%rd210+24];
ld.u64 %rd30, [%rd210+32];
ld.u64 %rd31, [%rd210+40];
// begin inline asm
{
.reg .u64 c;
.reg .u64 nc;
.reg .u64 t;
mad.lo.cc.u64 %rd15, %rd26, %rd27, 0;
madc.hi.cc.u64 c, %rd26, %rd27, 0;
madc.lo.cc.u64 %rd16, %rd26, %rd28, c;
madc.hi.cc.u64 c, %rd26, %rd28, 0;
madc.lo.cc.u64 %rd17, %rd26, %rd29, c;
madc.hi.cc.u64 c, %rd26, %rd29, 0;
madc.lo.cc.u64 %rd18, %rd26, %rd30, c;
madc.hi.cc.u64 c, %rd26, %rd30, 0;
madc.lo.cc.u64 %rd19, %rd26, %rd31, c;
madc.hi.u64 %rd20, %rd26, %rd31, 0;
mad.lo.cc.u64 %rd17, %rd27, %rd28, %rd17;
madc.hi.cc.u64 c, %rd27, %rd28, 0;
addc.cc.u64 t, %rd18, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd18, %rd27, %rd29, t;
madc.hi.cc.u64 c, %rd27, %rd29, nc;
addc.cc.u64 t, %rd19, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd19, %rd27, %rd30, t;
madc.hi.cc.u64 c, %rd27, %rd30, nc;
addc.cc.u64 t, %rd20, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd20, %rd27, %rd31, t;
madc.hi.u64 %rd21, %rd27, %rd31, nc;
mad.lo.cc.u64 %rd19, %rd28, %rd29, %rd19;
madc.hi.cc.u64 c, %rd28, %rd29, 0;
addc.cc.u64 t, %rd20, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd20, %rd28, %rd30, t;
madc.hi.cc.u64 c, %rd28, %rd30, nc;
addc.cc.u64 t, %rd21, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd21, %rd28, %rd31, t;
madc.hi.u64 %rd22, %rd28, %rd31, nc;
mad.lo.cc.u64 %rd21, %rd29, %rd30, %rd21;
madc.hi.cc.u64 c, %rd29, %rd30, 0;
addc.cc.u64 t, %rd22, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd22, %rd29, %rd31, t;
madc.hi.u64 %rd23, %rd29, %rd31, nc;
mad.lo.cc.u64 %rd23, %rd30, %rd31, %rd23;
madc.hi.u64 %rd24, %rd30, %rd31, 0;
}
// end inline asm
shr.u64 %rd194, %rd24, 63;
shl.b64 %rd224, %rd24, 1;
shr.u64 %rd225, %rd23, 63;
or.b64 %rd171, %rd224, %rd225;
shl.b64 %rd226, %rd23, 1;
shr.u64 %rd227, %rd22, 63;
or.b64 %rd148, %rd226, %rd227;
shl.b64 %rd228, %rd22, 1;
shr.u64 %rd229, %rd21, 63;
or.b64 %rd125, %rd228, %rd229;
shl.b64 %rd230, %rd21, 1;
shr.u64 %rd231, %rd20, 63;
or.b64 %rd102, %rd230, %rd231;
shl.b64 %rd232, %rd20, 1;
shr.u64 %rd233, %rd19, 63;
or.b64 %rd80, %rd232, %rd233;
shl.b64 %rd234, %rd19, 1;
shr.u64 %rd235, %rd18, 63;
or.b64 %rd79, %rd234, %rd235;
shl.b64 %rd236, %rd18, 1;
shr.u64 %rd237, %rd17, 63;
or.b64 %rd78, %rd236, %rd237;
shl.b64 %rd238, %rd17, 1;
shr.u64 %rd239, %rd16, 63;
or.b64 %rd77, %rd238, %rd239;
shl.b64 %rd240, %rd16, 1;
shr.u64 %rd241, %rd15, 63;
or.b64 %rd76, %rd240, %rd241;
shl.b64 %rd75, %rd15, 1;
// begin inline asm
{
mad.lo.cc.u64 %rd74, %rd26, %rd26, 0;
madc.hi.cc.u64 %rd75, %rd26, %rd26, %rd75;
madc.lo.cc.u64 %rd76, %rd27, %rd27, %rd76;
madc.hi.cc.u64 %rd77, %rd27, %rd27, %rd77;
madc.lo.cc.u64 %rd78, %rd28, %rd28, %rd78;
madc.hi.cc.u64 %rd79, %rd28, %rd28, %rd79;
madc.lo.cc.u64 %rd80, %rd29, %rd29, %rd80;
madc.hi.cc.u64 %rd102, %rd29, %rd29, %rd102;
madc.lo.cc.u64 %rd125, %rd30, %rd30, %rd125;
madc.hi.cc.u64 %rd148, %rd30, %rd30, %rd148;
madc.lo.cc.u64 %rd171, %rd31, %rd31, %rd171;
madc.hi.u64 %rd194, %rd31, %rd31, %rd194;
}
// end inline asm
mul.lo.s64 %rd88, %rd74, %rd211;
ld.u64 %rd82, [%rd13];
ld.u64 %rd83, [%rd13+8];
ld.u64 %rd84, [%rd13+16];
ld.u64 %rd85, [%rd13+24];
ld.u64 %rd86, [%rd13+32];
ld.u64 %rd87, [%rd13+40];
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd88, %rd82, %rd74;
madc.hi.cc.u64 c, %rd88, %rd82, 0;
addc.cc.u64 t, %rd75, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd75, %rd88, %rd83, t;
madc.hi.cc.u64 c, %rd88, %rd83, nc;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd88, %rd84, t;
madc.hi.cc.u64 c, %rd88, %rd84, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd88, %rd85, t;
madc.hi.cc.u64 c, %rd88, %rd85, nc;
addc.cc.u64 t, %rd78, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd78, %rd88, %rd86, t;
madc.hi.cc.u64 c, %rd88, %rd86, nc;
addc.cc.u64 t, %rd79, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd79, %rd88, %rd87, t;
madc.hi.cc.u64 c, %rd88, %rd87, nc;
addc.cc.u64 %rd80, %rd80, c;
addc.u64 %rd126, 0, 0;
}
// end inline asm
mul.lo.s64 %rd110, %rd75, %rd211;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd110, %rd82, %rd75;
madc.hi.cc.u64 c, %rd110, %rd82, 0;
addc.cc.u64 t, %rd76, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd76, %rd110, %rd83, t;
madc.hi.cc.u64 c, %rd110, %rd83, nc;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd110, %rd84, t;
madc.hi.cc.u64 c, %rd110, %rd84, nc;
addc.cc.u64 t, %rd78, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd78, %rd110, %rd85, t;
madc.hi.cc.u64 c, %rd110, %rd85, nc;
addc.cc.u64 t, %rd79, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd79, %rd110, %rd86, t;
madc.hi.cc.u64 c, %rd110, %rd86, nc;
addc.cc.u64 t, %rd80, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd80, %rd110, %rd87, t;
madc.hi.cc.u64 c, %rd110, %rd87, nc;
addc.cc.u64 c, c, %rd126;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd102, %rd102, c;
addc.u64 %rd126, nc, 0;
}
// end inline asm
mul.lo.s64 %rd133, %rd76, %rd211;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd133, %rd82, %rd76;
madc.hi.cc.u64 c, %rd133, %rd82, 0;
addc.cc.u64 t, %rd77, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd77, %rd133, %rd83, t;
madc.hi.cc.u64 c, %rd133, %rd83, nc;
addc.cc.u64 t, %rd78, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd78, %rd133, %rd84, t;
madc.hi.cc.u64 c, %rd133, %rd84, nc;
addc.cc.u64 t, %rd79, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd79, %rd133, %rd85, t;
madc.hi.cc.u64 c, %rd133, %rd85, nc;
addc.cc.u64 t, %rd80, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd80, %rd133, %rd86, t;
madc.hi.cc.u64 c, %rd133, %rd86, nc;
addc.cc.u64 t, %rd102, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd102, %rd133, %rd87, t;
madc.hi.cc.u64 c, %rd133, %rd87, nc;
addc.cc.u64 c, c, %rd126;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd125, %rd125, c;
addc.u64 %rd126, nc, 0;
}
// end inline asm
mul.lo.s64 %rd156, %rd77, %rd211;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd156, %rd82, %rd77;
madc.hi.cc.u64 c, %rd156, %rd82, 0;
addc.cc.u64 t, %rd78, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd78, %rd156, %rd83, t;
madc.hi.cc.u64 c, %rd156, %rd83, nc;
addc.cc.u64 t, %rd79, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd79, %rd156, %rd84, t;
madc.hi.cc.u64 c, %rd156, %rd84, nc;
addc.cc.u64 t, %rd80, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd80, %rd156, %rd85, t;
madc.hi.cc.u64 c, %rd156, %rd85, nc;
addc.cc.u64 t, %rd102, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd102, %rd156, %rd86, t;
madc.hi.cc.u64 c, %rd156, %rd86, nc;
addc.cc.u64 t, %rd125, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd125, %rd156, %rd87, t;
madc.hi.cc.u64 c, %rd156, %rd87, nc;
addc.cc.u64 c, c, %rd126;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd148, %rd148, c;
addc.u64 %rd126, nc, 0;
}
// end inline asm
mul.lo.s64 %rd179, %rd78, %rd211;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd179, %rd82, %rd78;
madc.hi.cc.u64 c, %rd179, %rd82, 0;
addc.cc.u64 t, %rd79, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd79, %rd179, %rd83, t;
madc.hi.cc.u64 c, %rd179, %rd83, nc;
addc.cc.u64 t, %rd80, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd80, %rd179, %rd84, t;
madc.hi.cc.u64 c, %rd179, %rd84, nc;
addc.cc.u64 t, %rd102, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd102, %rd179, %rd85, t;
madc.hi.cc.u64 c, %rd179, %rd85, nc;
addc.cc.u64 t, %rd125, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd125, %rd179, %rd86, t;
madc.hi.cc.u64 c, %rd179, %rd86, nc;
addc.cc.u64 t, %rd148, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd148, %rd179, %rd87, t;
madc.hi.cc.u64 c, %rd179, %rd87, nc;
addc.cc.u64 c, c, %rd126;
addc.u64 nc, 0, 0;
addc.cc.u64 %rd171, %rd171, c;
addc.u64 %rd126, nc, 0;
}
// end inline asm
mul.lo.s64 %rd202, %rd79, %rd211;
// begin inline asm
{
.reg .u64 c;
.reg .u64 t;
.reg .u64 nc;
mad.lo.cc.u64 c, %rd202, %rd82, %rd79;
madc.hi.cc.u64 c, %rd202, %rd82, 0;
addc.cc.u64 t, %rd80, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd80, %rd202, %rd83, t;
madc.hi.cc.u64 c, %rd202, %rd83, nc;
addc.cc.u64 t, %rd102, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd102, %rd202, %rd84, t;
madc.hi.cc.u64 c, %rd202, %rd84, nc;
addc.cc.u64 t, %rd125, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd125, %rd202, %rd85, t;
madc.hi.cc.u64 c, %rd202, %rd85, nc;
addc.cc.u64 t, %rd148, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd148, %rd202, %rd86, t;
madc.hi.cc.u64 c, %rd202, %rd86, nc;
addc.cc.u64 t, %rd171, c;
addc.u64 nc, 0, 0;
mad.lo.cc.u64 %rd171, %rd202, %rd87, t;
madc.hi.cc.u64 c, %rd202, %rd87, nc;
addc.cc.u64 c, c, %rd126;
add.u64 %rd194, %rd194, c;
}
// end inline asm
st.u8 [%rd12], %rd80;
shr.u64 %rd242, %rd80, 56;
st.u8 [%rd12+7], %rd242;
shr.u64 %rd243, %rd80, 48;
st.u8 [%rd12+6], %rd243;
shr.u64 %rd244, %rd80, 40;
st.u8 [%rd12+5], %rd244;
shr.u64 %rd245, %rd80, 32;
st.u8 [%rd12+4], %rd245;
shr.u64 %rd246, %rd80, 24;
st.u8 [%rd12+3], %rd246;
shr.u64 %rd247, %rd80, 16;
st.u8 [%rd12+2], %rd247;
shr.u64 %rd248, %rd80, 8;
st.u8 [%rd12+1], %rd248;
st.u8 [%rd12+8], %rd102;
shr.u64 %rd249, %rd102, 56;
st.u8 [%rd12+15], %rd249;
shr.u64 %rd250, %rd102, 48;
st.u8 [%rd12+14], %rd250;
shr.u64 %rd251, %rd102, 40;
st.u8 [%rd12+13], %rd251;
shr.u64 %rd252, %rd102, 32;
st.u8 [%rd12+12], %rd252;
shr.u64 %rd253, %rd102, 24;
st.u8 [%rd12+11], %rd253;
shr.u64 %rd254, %rd102, 16;
st.u8 [%rd12+10], %rd254;
shr.u64 %rd255, %rd102, 8;
st.u8 [%rd12+9], %rd255;
st.u8 [%rd12+16], %rd125;
shr.u64 %rd256, %rd125, 56;
st.u8 [%rd12+23], %rd256;
shr.u64 %rd257, %rd125, 48;
st.u8 [%rd12+22], %rd257;
shr.u64 %rd258, %rd125, 40;
st.u8 [%rd12+21], %rd258;
shr.u64 %rd259, %rd125, 32;
st.u8 [%rd12+20], %rd259;
shr.u64 %rd260, %rd125, 24;
st.u8 [%rd12+19], %rd260;
shr.u64 %rd261, %rd125, 16;
st.u8 [%rd12+18], %rd261;
shr.u64 %rd262, %rd125, 8;
st.u8 [%rd12+17], %rd262;
st.u8 [%rd12+24], %rd148;
shr.u64 %rd263, %rd148, 56;
st.u8 [%rd12+31], %rd263;
shr.u64 %rd264, %rd148, 48;
st.u8 [%rd12+30], %rd264;
shr.u64 %rd265, %rd148, 40;
st.u8 [%rd12+29], %rd265;
shr.u64 %rd266, %rd148, 32;
st.u8 [%rd12+28], %rd266;
shr.u64 %rd267, %rd148, 24;
st.u8 [%rd12+27], %rd267;
shr.u64 %rd268, %rd148, 16;
st.u8 [%rd12+26], %rd268;
shr.u64 %rd269, %rd148, 8;
st.u8 [%rd12+25], %rd269;
st.u8 [%rd12+32], %rd171;
shr.u64 %rd270, %rd171, 56;
st.u8 [%rd12+39], %rd270;
shr.u64 %rd271, %rd171, 48;
st.u8 [%rd12+38], %rd271;
shr.u64 %rd272, %rd171, 40;
st.u8 [%rd12+37], %rd272;
shr.u64 %rd273, %rd171, 32;
st.u8 [%rd12+36], %rd273;
shr.u64 %rd274, %rd171, 24;
st.u8 [%rd12+35], %rd274;
shr.u64 %rd275, %rd171, 16;
st.u8 [%rd12+34], %rd275;
shr.u64 %rd276, %rd171, 8;
st.u8 [%rd12+33], %rd276;
st.u8 [%rd12+40], %rd194;
shr.u64 %rd277, %rd194, 56;
st.u8 [%rd12+47], %rd277;
shr.u64 %rd278, %rd194, 48;
st.u8 [%rd12+46], %rd278;
shr.u64 %rd279, %rd194, 40;
st.u8 [%rd12+45], %rd279;
shr.u64 %rd280, %rd194, 32;
st.u8 [%rd12+44], %rd280;
shr.u64 %rd281, %rd194, 24;
st.u8 [%rd12+43], %rd281;
shr.u64 %rd282, %rd194, 16;
st.u8 [%rd12+42], %rd282;
shr.u64 %rd283, %rd194, 8;
st.u8 [%rd12+41], %rd283;
ld.u64 %rd7, [%rd13+40];
setp.lt.u64 %p1, %rd194, %rd7;
@%p1 bra $L__BB1_12;
setp.gt.u64 %p2, %rd194, %rd7;
ld.u64 %rd8, [%rd13+32];
@%p2 bra $L__BB1_11;
setp.lt.u64 %p3, %rd171, %rd8;
@%p3 bra $L__BB1_12;
setp.gt.u64 %p4, %rd171, %rd8;
@%p4 bra $L__BB1_11;
ld.u64 %rd9, [%rd13+24];
setp.lt.u64 %p5, %rd148, %rd9;
@%p5 bra $L__BB1_12;
setp.gt.u64 %p6, %rd148, %rd9;
@%p6 bra $L__BB1_11;
ld.u64 %rd10, [%rd13+16];
setp.lt.u64 %p7, %rd125, %rd10;
@%p7 bra $L__BB1_12;
setp.gt.u64 %p8, %rd125, %rd10;
@%p8 bra $L__BB1_11;
ld.u64 %rd11, [%rd13+8];
setp.lt.u64 %p9, %rd102, %rd11;
@%p9 bra $L__BB1_12;
setp.gt.u64 %p10, %rd102, %rd11;
@%p10 bra $L__BB1_11;
ld.u64 %rd284, [%rd13];
setp.lt.u64 %p11, %rd80, %rd284;
@%p11 bra $L__BB1_12;
$L__BB1_11:
ld.param.u64 %rd345, [_Z12sqr_mont_384PyPKyS1_y_param_0];
ld.u64 %rd297, [%rd13];
ld.u64 %rd298, [%rd13+8];
ld.u64 %rd299, [%rd13+16];
ld.u64 %rd300, [%rd13+24];
// begin inline asm
sub.cc.u64 %rd285, %rd80, %rd297;
subc.cc.u64 %rd286, %rd102, %rd298;
subc.cc.u64 %rd287, %rd125, %rd299;
subc.cc.u64 %rd288, %rd148, %rd300;
subc.cc.u64 %rd289, %rd171, %rd8;
subc.u64 %rd290, %rd194, %rd7;
// end inline asm
st.u8 [%rd345], %rd285;
shr.u64 %rd303, %rd285, 56;
st.u8 [%rd345+7], %rd303;
shr.u64 %rd304, %rd285, 48;
st.u8 [%rd345+6], %rd304;
shr.u64 %rd305, %rd285, 40;
st.u8 [%rd345+5], %rd305;
shr.u64 %rd306, %rd285, 32;
st.u8 [%rd345+4], %rd306;
shr.u64 %rd307, %rd285, 24;
st.u8 [%rd345+3], %rd307;
shr.u64 %rd308, %rd285, 16;
st.u8 [%rd345+2], %rd308;
shr.u64 %rd309, %rd285, 8;
st.u8 [%rd345+1], %rd309;
st.u8 [%rd345+8], %rd286;
shr.u64 %rd310, %rd286, 56;
st.u8 [%rd345+15], %rd310;
shr.u64 %rd311, %rd286, 48;
st.u8 [%rd345+14], %rd311;
shr.u64 %rd312, %rd286, 40;
st.u8 [%rd345+13], %rd312;
shr.u64 %rd313, %rd286, 32;
st.u8 [%rd345+12], %rd313;
shr.u64 %rd314, %rd286, 24;
st.u8 [%rd345+11], %rd314;
shr.u64 %rd315, %rd286, 16;
st.u8 [%rd345+10], %rd315;
shr.u64 %rd316, %rd286, 8;
st.u8 [%rd345+9], %rd316;
st.u8 [%rd345+16], %rd287;
shr.u64 %rd317, %rd287, 56;
st.u8 [%rd345+23], %rd317;
shr.u64 %rd318, %rd287, 48;
st.u8 [%rd345+22], %rd318;
shr.u64 %rd319, %rd287, 40;
st.u8 [%rd345+21], %rd319;
shr.u64 %rd320, %rd287, 32;
st.u8 [%rd345+20], %rd320;
shr.u64 %rd321, %rd287, 24;
st.u8 [%rd345+19], %rd321;
shr.u64 %rd322, %rd287, 16;
st.u8 [%rd345+18], %rd322;
shr.u64 %rd323, %rd287, 8;
st.u8 [%rd345+17], %rd323;
st.u8 [%rd345+24], %rd288;
shr.u64 %rd324, %rd288, 56;
st.u8 [%rd345+31], %rd324;
shr.u64 %rd325, %rd288, 48;
st.u8 [%rd345+30], %rd325;
shr.u64 %rd326, %rd288, 40;
st.u8 [%rd345+29], %rd326;
shr.u64 %rd327, %rd288, 32;
st.u8 [%rd345+28], %rd327;
shr.u64 %rd328, %rd288, 24;
st.u8 [%rd345+27], %rd328;
shr.u64 %rd329, %rd288, 16;
st.u8 [%rd345+26], %rd329;
shr.u64 %rd330, %rd288, 8;
st.u8 [%rd345+25], %rd330;
st.u8 [%rd345+32], %rd289;
shr.u64 %rd331, %rd289, 56;
st.u8 [%rd345+39], %rd331;
shr.u64 %rd332, %rd289, 48;
st.u8 [%rd345+38], %rd332;
shr.u64 %rd333, %rd289, 40;
st.u8 [%rd345+37], %rd333;
shr.u64 %rd334, %rd289, 32;
st.u8 [%rd345+36], %rd334;
shr.u64 %rd335, %rd289, 24;
st.u8 [%rd345+35], %rd335;
shr.u64 %rd336, %rd289, 16;
st.u8 [%rd345+34], %rd336;
shr.u64 %rd337, %rd289, 8;
st.u8 [%rd345+33], %rd337;
st.u8 [%rd345+40], %rd290;
shr.u64 %rd338, %rd290, 56;
st.u8 [%rd345+47], %rd338;
shr.u64 %rd339, %rd290, 48;
st.u8 [%rd345+46], %rd339;
shr.u64 %rd340, %rd290, 40;
st.u8 [%rd345+45], %rd340;
shr.u64 %rd341, %rd290, 32;
st.u8 [%rd345+44], %rd341;
shr.u64 %rd342, %rd290, 24;
st.u8 [%rd345+43], %rd342;
shr.u64 %rd343, %rd290, 16;
st.u8 [%rd345+42], %rd343;
shr.u64 %rd344, %rd290, 8;
st.u8 [%rd345+41], %rd344;
$L__BB1_12:
ret;
}
// .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 .pred %p<12>;
.reg .b64 %rd<96>;
ld.param.u64 %rd13, [_Z11add_mod_384PyPKyS1_S1__param_0];
ld.param.u64 %rd33, [_Z11add_mod_384PyPKyS1_S1__param_1];
ld.param.u64 %rd34, [_Z11add_mod_384PyPKyS1_S1__param_2];
ld.param.u64 %rd14, [_Z11add_mod_384PyPKyS1_S1__param_3];
ld.u64 %rd21, [%rd33];
ld.u64 %rd22, [%rd33+8];
ld.u64 %rd23, [%rd33+16];
ld.u64 %rd24, [%rd33+24];
ld.u64 %rd25, [%rd33+32];
ld.u64 %rd26, [%rd33+40];
ld.u64 %rd27, [%rd34];
ld.u64 %rd28, [%rd34+8];
ld.u64 %rd29, [%rd34+16];
ld.u64 %rd30, [%rd34+24];
ld.u64 %rd31, [%rd34+32];
ld.u64 %rd32, [%rd34+40];
// begin inline asm
add.cc.u64 %rd15, %rd21, %rd27;
addc.cc.u64 %rd16, %rd22, %rd28;
addc.cc.u64 %rd17, %rd23, %rd29;
addc.cc.u64 %rd18, %rd24, %rd30;
addc.cc.u64 %rd19, %rd25, %rd31;
addc.u64 %rd20, %rd26, %rd32;
// end inline asm
st.u64 [%rd13], %rd15;
st.u64 [%rd13+8], %rd16;
st.u64 [%rd13+16], %rd17;
st.u64 [%rd13+24], %rd18;
st.u64 [%rd13+32], %rd19;
st.u64 [%rd13+40], %rd20;
ld.u64 %rd8, [%rd14+40];
setp.lt.u64 %p1, %rd20, %rd8;
@%p1 bra $L__BB2_12;
setp.gt.u64 %p2, %rd20, %rd8;
ld.u64 %rd9, [%rd14+32];
@%p2 bra $L__BB2_11;
setp.lt.u64 %p3, %rd19, %rd9;
@%p3 bra $L__BB2_12;
setp.gt.u64 %p4, %rd19, %rd9;
@%p4 bra $L__BB2_11;
ld.u64 %rd10, [%rd14+24];
setp.lt.u64 %p5, %rd18, %rd10;
@%p5 bra $L__BB2_12;
setp.gt.u64 %p6, %rd18, %rd10;
@%p6 bra $L__BB2_11;
ld.u64 %rd11, [%rd14+16];
setp.lt.u64 %p7, %rd17, %rd11;
@%p7 bra $L__BB2_12;
setp.gt.u64 %p8, %rd17, %rd11;
@%p8 bra $L__BB2_11;
ld.u64 %rd12, [%rd14+8];
setp.lt.u64 %p9, %rd16, %rd12;
@%p9 bra $L__BB2_12;
setp.gt.u64 %p10, %rd16, %rd12;
@%p10 bra $L__BB2_11;
ld.u64 %rd35, [%rd14];
setp.lt.u64 %p11, %rd15, %rd35;
@%p11 bra $L__BB2_12;
$L__BB2_11:
ld.u64 %rd48, [%rd14];
ld.u64 %rd49, [%rd14+8];
ld.u64 %rd50, [%rd14+16];
ld.u64 %rd51, [%rd14+24];
// begin inline asm
sub.cc.u64 %rd36, %rd15, %rd48;
subc.cc.u64 %rd37, %rd16, %rd49;
subc.cc.u64 %rd38, %rd17, %rd50;
subc.cc.u64 %rd39, %rd18, %rd51;
subc.cc.u64 %rd40, %rd19, %rd9;
subc.u64 %rd41, %rd20, %rd8;
// end inline asm
st.u8 [%rd13], %rd36;
shr.u64 %rd54, %rd36, 56;
st.u8 [%rd13+7], %rd54;
shr.u64 %rd55, %rd36, 48;
st.u8 [%rd13+6], %rd55;
shr.u64 %rd56, %rd36, 40;
st.u8 [%rd13+5], %rd56;
shr.u64 %rd57, %rd36, 32;
st.u8 [%rd13+4], %rd57;
shr.u64 %rd58, %rd36, 24;
st.u8 [%rd13+3], %rd58;
shr.u64 %rd59, %rd36, 16;
st.u8 [%rd13+2], %rd59;
shr.u64 %rd60, %rd36, 8;
st.u8 [%rd13+1], %rd60;
st.u8 [%rd13+8], %rd37;
shr.u64 %rd61, %rd37, 56;
st.u8 [%rd13+15], %rd61;
shr.u64 %rd62, %rd37, 48;
st.u8 [%rd13+14], %rd62;
shr.u64 %rd63, %rd37, 40;
st.u8 [%rd13+13], %rd63;
shr.u64 %rd64, %rd37, 32;
st.u8 [%rd13+12], %rd64;
shr.u64 %rd65, %rd37, 24;
st.u8 [%rd13+11], %rd65;
shr.u64 %rd66, %rd37, 16;
st.u8 [%rd13+10], %rd66;
shr.u64 %rd67, %rd37, 8;
st.u8 [%rd13+9], %rd67;
st.u8 [%rd13+16], %rd38;
shr.u64 %rd68, %rd38, 56;
st.u8 [%rd13+23], %rd68;
shr.u64 %rd69, %rd38, 48;
st.u8 [%rd13+22], %rd69;
shr.u64 %rd70, %rd38, 40;
st.u8 [%rd13+21], %rd70;
shr.u64 %rd71, %rd38, 32;
st.u8 [%rd13+20], %rd71;
shr.u64 %rd72, %rd38, 24;
st.u8 [%rd13+19], %rd72;
shr.u64 %rd73, %rd38, 16;
st.u8 [%rd13+18], %rd73;
shr.u64 %rd74, %rd38, 8;
st.u8 [%rd13+17], %rd74;
st.u8 [%rd13+24], %rd39;
shr.u64 %rd75, %rd39, 56;
st.u8 [%rd13+31], %rd75;
shr.u64 %rd76, %rd39, 48;
st.u8 [%rd13+30], %rd76;
shr.u64 %rd77, %rd39, 40;
st.u8 [%rd13+29], %rd77;
shr.u64 %rd78, %rd39, 32;
st.u8 [%rd13+28], %rd78;
shr.u64 %rd79, %rd39, 24;
st.u8 [%rd13+27], %rd79;
shr.u64 %rd80, %rd39, 16;
st.u8 [%rd13+26], %rd80;
shr.u64 %rd81, %rd39, 8;
st.u8 [%rd13+25], %rd81;
st.u8 [%rd13+32], %rd40;
shr.u64 %rd82, %rd40, 56;
st.u8 [%rd13+39], %rd82;
shr.u64 %rd83, %rd40, 48;
st.u8 [%rd13+38], %rd83;
shr.u64 %rd84, %rd40, 40;
st.u8 [%rd13+37], %rd84;
shr.u64 %rd85, %rd40, 32;
st.u8 [%rd13+36], %rd85;
shr.u64 %rd86, %rd40, 24;
st.u8 [%rd13+35], %rd86;
shr.u64 %rd87, %rd40, 16;
st.u8 [%rd13+34], %rd87;
shr.u64 %rd88, %rd40, 8;
st.u8 [%rd13+33], %rd88;
st.u8 [%rd13+40], %rd41;
shr.u64 %rd89, %rd41, 56;
st.u8 [%rd13+47], %rd89;
shr.u64 %rd90, %rd41, 48;
st.u8 [%rd13+46], %rd90;
shr.u64 %rd91, %rd41, 40;
st.u8 [%rd13+45], %rd91;
shr.u64 %rd92, %rd41, 32;
st.u8 [%rd13+44], %rd92;
shr.u64 %rd93, %rd41, 24;
st.u8 [%rd13+43], %rd93;
shr.u64 %rd94, %rd41, 16;
st.u8 [%rd13+42], %rd94;
shr.u64 %rd95, %rd41, 8;
st.u8 [%rd13+41], %rd95;
$L__BB2_12:
ret;
}
// .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
)
{
.reg .pred %p<12>;
.reg .b64 %rd<155>;
ld.param.u64 %rd25, [_Z11sub_mod_384PyPKyS1_S1__param_0];
ld.param.u64 %rd27, [_Z11sub_mod_384PyPKyS1_S1__param_1];
ld.param.u64 %rd6, [_Z11sub_mod_384PyPKyS1_S1__param_2];
ld.param.u64 %rd26, [_Z11sub_mod_384PyPKyS1_S1__param_3];
ld.u8 %rd28, [%rd27];
ld.u8 %rd29, [%rd27+1];
bfi.b64 %rd30, %rd29, %rd28, 8, 8;
ld.u8 %rd31, [%rd27+2];
ld.u8 %rd32, [%rd27+3];
bfi.b64 %rd33, %rd32, %rd31, 8, 8;
bfi.b64 %rd34, %rd33, %rd30, 16, 16;
ld.u8 %rd35, [%rd27+4];
ld.u8 %rd36, [%rd27+5];
bfi.b64 %rd37, %rd36, %rd35, 8, 8;
ld.u8 %rd38, [%rd27+6];
ld.u8 %rd39, [%rd27+7];
bfi.b64 %rd40, %rd39, %rd38, 8, 8;
bfi.b64 %rd41, %rd40, %rd37, 16, 16;
bfi.b64 %rd154, %rd41, %rd34, 32, 32;
ld.u8 %rd42, [%rd27+8];
ld.u8 %rd43, [%rd27+9];
bfi.b64 %rd44, %rd43, %rd42, 8, 8;
ld.u8 %rd45, [%rd27+10];
ld.u8 %rd46, [%rd27+11];
bfi.b64 %rd47, %rd46, %rd45, 8, 8;
bfi.b64 %rd48, %rd47, %rd44, 16, 16;
ld.u8 %rd49, [%rd27+12];
ld.u8 %rd50, [%rd27+13];
bfi.b64 %rd51, %rd50, %rd49, 8, 8;
ld.u8 %rd52, [%rd27+14];
ld.u8 %rd53, [%rd27+15];
bfi.b64 %rd54, %rd53, %rd52, 8, 8;
bfi.b64 %rd55, %rd54, %rd51, 16, 16;
bfi.b64 %rd153, %rd55, %rd48, 32, 32;
ld.u8 %rd56, [%rd27+16];
ld.u8 %rd57, [%rd27+17];
bfi.b64 %rd58, %rd57, %rd56, 8, 8;
ld.u8 %rd59, [%rd27+18];
ld.u8 %rd60, [%rd27+19];
bfi.b64 %rd61, %rd60, %rd59, 8, 8;
bfi.b64 %rd62, %rd61, %rd58, 16, 16;
ld.u8 %rd63, [%rd27+20];
ld.u8 %rd64, [%rd27+21];
bfi.b64 %rd65, %rd64, %rd63, 8, 8;
ld.u8 %rd66, [%rd27+22];
ld.u8 %rd67, [%rd27+23];
bfi.b64 %rd68, %rd67, %rd66, 8, 8;
bfi.b64 %rd69, %rd68, %rd65, 16, 16;
bfi.b64 %rd152, %rd69, %rd62, 32, 32;
ld.u8 %rd70, [%rd27+24];
ld.u8 %rd71, [%rd27+25];
bfi.b64 %rd72, %rd71, %rd70, 8, 8;
ld.u8 %rd73, [%rd27+26];
ld.u8 %rd74, [%rd27+27];
bfi.b64 %rd75, %rd74, %rd73, 8, 8;
bfi.b64 %rd76, %rd75, %rd72, 16, 16;
ld.u8 %rd77, [%rd27+28];
ld.u8 %rd78, [%rd27+29];
bfi.b64 %rd79, %rd78, %rd77, 8, 8;
ld.u8 %rd80, [%rd27+30];
ld.u8 %rd81, [%rd27+31];
bfi.b64 %rd82, %rd81, %rd80, 8, 8;
bfi.b64 %rd83, %rd82, %rd79, 16, 16;
bfi.b64 %rd151, %rd83, %rd76, 32, 32;
ld.u8 %rd84, [%rd27+32];
ld.u8 %rd85, [%rd27+33];
bfi.b64 %rd86, %rd85, %rd84, 8, 8;
ld.u8 %rd87, [%rd27+34];
ld.u8 %rd88, [%rd27+35];
bfi.b64 %rd89, %rd88, %rd87, 8, 8;
bfi.b64 %rd90, %rd89, %rd86, 16, 16;
ld.u8 %rd91, [%rd27+36];
ld.u8 %rd92, [%rd27+37];
bfi.b64 %rd93, %rd92, %rd91, 8, 8;
ld.u8 %rd94, [%rd27+38];
ld.u8 %rd95, [%rd27+39];
bfi.b64 %rd96, %rd95, %rd94, 8, 8;
bfi.b64 %rd97, %rd96, %rd93, 16, 16;
bfi.b64 %rd150, %rd97, %rd90, 32, 32;
ld.u64 %rd7, [%rd6+40];
ld.u8 %rd98, [%rd27+40];
ld.u8 %rd99, [%rd27+41];
bfi.b64 %rd100, %rd99, %rd98, 8, 8;
ld.u8 %rd101, [%rd27+42];
ld.u8 %rd102, [%rd27+43];
bfi.b64 %rd103, %rd102, %rd101, 8, 8;
bfi.b64 %rd104, %rd103, %rd100, 16, 16;
ld.u8 %rd105, [%rd27+44];
ld.u8 %rd106, [%rd27+45];
bfi.b64 %rd107, %rd106, %rd105, 8, 8;
ld.u8 %rd108, [%rd27+46];
ld.u8 %rd109, [%rd27+47];
bfi.b64 %rd110, %rd109, %rd108, 8, 8;
bfi.b64 %rd111, %rd110, %rd107, 16, 16;
bfi.b64 %rd149, %rd111, %rd104, 32, 32;
setp.lt.u64 %p1, %rd7, %rd149;
@%p1 bra $L__BB3_12;
setp.gt.u64 %p2, %rd7, %rd149;
@%p2 bra $L__BB3_11;
ld.u64 %rd9, [%rd6+32];
setp.lt.u64 %p3, %rd9, %rd150;
@%p3 bra $L__BB3_12;
setp.gt.u64 %p4, %rd9, %rd150;
@%p4 bra $L__BB3_11;
ld.u64 %rd10, [%rd6+24];
setp.lt.u64 %p5, %rd10, %rd151;
@%p5 bra $L__BB3_12;
setp.gt.u64 %p6, %rd10, %rd151;
@%p6 bra $L__BB3_11;
ld.u64 %rd11, [%rd6+16];
setp.lt.u64 %p7, %rd11, %rd152;
@%p7 bra $L__BB3_12;
setp.gt.u64 %p8, %rd11, %rd152;
@%p8 bra $L__BB3_11;
ld.u64 %rd12, [%rd6+8];
setp.lt.u64 %p9, %rd12, %rd153;
@%p9 bra $L__BB3_12;
setp.gt.u64 %p10, %rd12, %rd153;
@%p10 bra $L__BB3_11;
ld.u64 %rd112, [%rd6];
setp.le.u64 %p11, %rd112, %rd154;
@%p11 bra $L__BB3_12;
$L__BB3_11:
ld.u64 %rd125, [%rd26];
ld.u64 %rd126, [%rd26+8];
ld.u64 %rd127, [%rd26+16];
ld.u64 %rd128, [%rd26+24];
ld.u64 %rd129, [%rd26+32];
ld.u64 %rd130, [%rd26+40];
// begin inline asm
add.cc.u64 %rd154, %rd154, %rd125;
addc.cc.u64 %rd153, %rd153, %rd126;
addc.cc.u64 %rd152, %rd152, %rd127;
addc.cc.u64 %rd151, %rd151, %rd128;
addc.cc.u64 %rd150, %rd150, %rd129;
addc.u64 %rd149, %rd149, %rd130;
// end inline asm
$L__BB3_12:
ld.u64 %rd143, [%rd6];
ld.u64 %rd144, [%rd6+8];
ld.u64 %rd145, [%rd6+16];
ld.u64 %rd146, [%rd6+24];
ld.u64 %rd147, [%rd6+32];
// begin inline asm
sub.cc.u64 %rd131, %rd154, %rd143;
subc.cc.u64 %rd132, %rd153, %rd144;
subc.cc.u64 %rd133, %rd152, %rd145;
subc.cc.u64 %rd134, %rd151, %rd146;
subc.cc.u64 %rd135, %rd150, %rd147;
subc.u64 %rd136, %rd149, %rd7;
// end inline asm
st.u64 [%rd25], %rd131;
st.u64 [%rd25+8], %rd132;
st.u64 [%rd25+16], %rd133;
st.u64 [%rd25+24], %rd134;
st.u64 [%rd25+32], %rd135;
st.u64 [%rd25+40], %rd136;
ret;
}
// .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<22>;
ld.param.u64 %rd19, [_Z18sub_mod_384_unsafePyPKyS1__param_0];
ld.param.u64 %rd20, [_Z18sub_mod_384_unsafePyPKyS1__param_1];
ld.param.u64 %rd21, [_Z18sub_mod_384_unsafePyPKyS1__param_2];
ld.u64 %rd7, [%rd20];
ld.u64 %rd8, [%rd20+8];
ld.u64 %rd9, [%rd20+16];
ld.u64 %rd10, [%rd20+24];
ld.u64 %rd11, [%rd20+32];
ld.u64 %rd12, [%rd20+40];
ld.u64 %rd13, [%rd21];
ld.u64 %rd14, [%rd21+8];
ld.u64 %rd15, [%rd21+16];
ld.u64 %rd16, [%rd21+24];
ld.u64 %rd17, [%rd21+32];
ld.u64 %rd18, [%rd21+40];
// begin inline asm
sub.cc.u64 %rd1, %rd7, %rd13;
subc.cc.u64 %rd2, %rd8, %rd14;
subc.cc.u64 %rd3, %rd9, %rd15;
subc.cc.u64 %rd4, %rd10, %rd16;
subc.cc.u64 %rd5, %rd11, %rd17;
subc.u64 %rd6, %rd12, %rd18;
// end inline asm
st.u64 [%rd19], %rd1;
st.u64 [%rd19+8], %rd2;
st.u64 [%rd19+16], %rd3;
st.u64 [%rd19+24], %rd4;
st.u64 [%rd19+32], %rd5;
st.u64 [%rd19+40], %rd6;
ret;
}
// .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<22>;
ld.param.u64 %rd19, [_Z18add_mod_384_unsafePyPKyS1__param_0];
ld.param.u64 %rd20, [_Z18add_mod_384_unsafePyPKyS1__param_1];
ld.param.u64 %rd21, [_Z18add_mod_384_unsafePyPKyS1__param_2];
ld.u64 %rd7, [%rd20];
ld.u64 %rd8, [%rd20+8];
ld.u64 %rd9, [%rd20+16];
ld.u64 %rd10, [%rd20+24];
ld.u64 %rd11, [%rd20+32];
ld.u64 %rd12, [%rd20+40];
ld.u64 %rd13, [%rd21];
ld.u64 %rd14, [%rd21+8];
ld.u64 %rd15, [%rd21+16];
ld.u64 %rd16, [%rd21+24];
ld.u64 %rd17, [%rd21+32];
ld.u64 %rd18, [%rd21+40];
// begin inline asm
add.cc.u64 %rd1, %rd7, %rd13;
addc.cc.u64 %rd2, %rd8, %rd14;
addc.cc.u64 %rd3, %rd9, %rd15;
addc.cc.u64 %rd4, %rd10, %rd16;
addc.cc.u64 %rd5, %rd11, %rd17;
addc.u64 %rd6, %rd12, %rd18;
// end inline asm
st.u64 [%rd19], %rd1;
st.u64 [%rd19+8], %rd2;
st.u64 [%rd19+16], %rd3;
st.u64 [%rd19+24], %rd4;
st.u64 [%rd19+32], %rd5;
st.u64 [%rd19+40], %rd6;
ret;
}
// .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<30>;
ld.param.u64 %rd1, [_Z16div_by_2_mod_384PyPKy_param_0];
ld.param.u64 %rd2, [_Z16div_by_2_mod_384PyPKy_param_1];
ld.u64 %rd3, [%rd2+8];
shl.b64 %rd4, %rd3, 63;
ld.u64 %rd5, [%rd2];
shr.u64 %rd6, %rd5, 1;
or.b64 %rd7, %rd6, %rd4;
st.u64 [%rd1], %rd7;
ld.u64 %rd8, [%rd2+16];
shl.b64 %rd9, %rd8, 63;
ld.u64 %rd10, [%rd2+8];
shr.u64 %rd11, %rd10, 1;
or.b64 %rd12, %rd11, %rd9;
st.u64 [%rd1+8], %rd12;
ld.u64 %rd13, [%rd2+24];
shl.b64 %rd14, %rd13, 63;
ld.u64 %rd15, [%rd2+16];
shr.u64 %rd16, %rd15, 1;
or.b64 %rd17, %rd16, %rd14;
st.u64 [%rd1+16], %rd17;
ld.u64 %rd18, [%rd2+32];
shl.b64 %rd19, %rd18, 63;
ld.u64 %rd20, [%rd2+24];
shr.u64 %rd21, %rd20, 1;
or.b64 %rd22, %rd21, %rd19;
st.u64 [%rd1+24], %rd22;
ld.u64 %rd23, [%rd2+40];
shl.b64 %rd24, %rd23, 63;
ld.u64 %rd25, [%rd2+32];
shr.u64 %rd26, %rd25, 1;
or.b64 %rd27, %rd26, %rd24;
st.u64 [%rd1+32], %rd27;
ld.u64 %rd28, [%rd2+40];
shr.u64 %rd29, %rd28, 1;
st.u64 [%rd1+40], %rd29;
ret;
}
// .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<14>;
.reg .b16 %rs<3>;
.reg .b32 %r<5>;
.reg .b64 %rd<161>;
ld.param.u64 %rd27, [_Z12cneg_mod_384PyPKybS1__param_0];
ld.param.s8 %rs1, [_Z12cneg_mod_384PyPKybS1__param_2];
ld.param.u64 %rd28, [_Z12cneg_mod_384PyPKybS1__param_1];
ld.param.u64 %rd29, [_Z12cneg_mod_384PyPKybS1__param_3];
setp.eq.s16 %p1, %rs1, 0;
@%p1 bra $L__BB7_14;
ld.u8 %rd30, [%rd29];
ld.u8 %rd31, [%rd29+1];
bfi.b64 %rd32, %rd31, %rd30, 8, 8;
ld.u8 %rd33, [%rd29+2];
ld.u8 %rd34, [%rd29+3];
bfi.b64 %rd35, %rd34, %rd33, 8, 8;
bfi.b64 %rd36, %rd35, %rd32, 16, 16;
ld.u8 %rd37, [%rd29+4];
ld.u8 %rd38, [%rd29+5];
bfi.b64 %rd39, %rd38, %rd37, 8, 8;
ld.u8 %rd40, [%rd29+6];
ld.u8 %rd41, [%rd29+7];
bfi.b64 %rd42, %rd41, %rd40, 8, 8;
bfi.b64 %rd43, %rd42, %rd39, 16, 16;
bfi.b64 %rd159, %rd43, %rd36, 32, 32;
ld.u8 %rd44, [%rd29+8];
ld.u8 %rd45, [%rd29+9];
bfi.b64 %rd46, %rd45, %rd44, 8, 8;
ld.u8 %rd47, [%rd29+10];
ld.u8 %rd48, [%rd29+11];
bfi.b64 %rd49, %rd48, %rd47, 8, 8;
bfi.b64 %rd50, %rd49, %rd46, 16, 16;
ld.u8 %rd51, [%rd29+12];
ld.u8 %rd52, [%rd29+13];
bfi.b64 %rd53, %rd52, %rd51, 8, 8;
ld.u8 %rd54, [%rd29+14];
ld.u8 %rd55, [%rd29+15];
bfi.b64 %rd56, %rd55, %rd54, 8, 8;
bfi.b64 %rd57, %rd56, %rd53, 16, 16;
bfi.b64 %rd158, %rd57, %rd50, 32, 32;
ld.u8 %rd58, [%rd29+16];
ld.u8 %rd59, [%rd29+17];
bfi.b64 %rd60, %rd59, %rd58, 8, 8;
ld.u8 %rd61, [%rd29+18];
ld.u8 %rd62, [%rd29+19];
bfi.b64 %rd63, %rd62, %rd61, 8, 8;
bfi.b64 %rd64, %rd63, %rd60, 16, 16;
ld.u8 %rd65, [%rd29+20];
ld.u8 %rd66, [%rd29+21];
bfi.b64 %rd67, %rd66, %rd65, 8, 8;
ld.u8 %rd68, [%rd29+22];
ld.u8 %rd69, [%rd29+23];
bfi.b64 %rd70, %rd69, %rd68, 8, 8;
bfi.b64 %rd71, %rd70, %rd67, 16, 16;
bfi.b64 %rd157, %rd71, %rd64, 32, 32;
ld.u8 %rd72, [%rd29+24];
ld.u8 %rd73, [%rd29+25];
bfi.b64 %rd74, %rd73, %rd72, 8, 8;
ld.u8 %rd75, [%rd29+26];
ld.u8 %rd76, [%rd29+27];
bfi.b64 %rd77, %rd76, %rd75, 8, 8;
bfi.b64 %rd78, %rd77, %rd74, 16, 16;
ld.u8 %rd79, [%rd29+28];
ld.u8 %rd80, [%rd29+29];
bfi.b64 %rd81, %rd80, %rd79, 8, 8;
ld.u8 %rd82, [%rd29+30];
ld.u8 %rd83, [%rd29+31];
bfi.b64 %rd84, %rd83, %rd82, 8, 8;
bfi.b64 %rd85, %rd84, %rd81, 16, 16;
bfi.b64 %rd156, %rd85, %rd78, 32, 32;
ld.u8 %rd86, [%rd29+32];
ld.u8 %rd87, [%rd29+33];
bfi.b64 %rd88, %rd87, %rd86, 8, 8;
ld.u8 %rd89, [%rd29+34];
ld.u8 %rd90, [%rd29+35];
bfi.b64 %rd91, %rd90, %rd89, 8, 8;
bfi.b64 %rd92, %rd91, %rd88, 16, 16;
ld.u8 %rd93, [%rd29+36];
ld.u8 %rd94, [%rd29+37];
bfi.b64 %rd95, %rd94, %rd93, 8, 8;
ld.u8 %rd96, [%rd29+38];
ld.u8 %rd97, [%rd29+39];
bfi.b64 %rd98, %rd97, %rd96, 8, 8;
bfi.b64 %rd99, %rd98, %rd95, 16, 16;
bfi.b64 %rd155, %rd99, %rd92, 32, 32;
ld.u64 %rd7, [%rd28+40];
ld.u8 %rd100, [%rd29+40];
ld.u8 %rd101, [%rd29+41];
bfi.b64 %rd102, %rd101, %rd100, 8, 8;
ld.u8 %rd103, [%rd29+42];
ld.u8 %rd104, [%rd29+43];
bfi.b64 %rd105, %rd104, %rd103, 8, 8;
bfi.b64 %rd106, %rd105, %rd102, 16, 16;
ld.u8 %rd107, [%rd29+44];
ld.u8 %rd108, [%rd29+45];
bfi.b64 %rd109, %rd108, %rd107, 8, 8;
ld.u8 %rd110, [%rd29+46];
ld.u8 %rd111, [%rd29+47];
bfi.b64 %rd112, %rd111, %rd110, 8, 8;
bfi.b64 %rd113, %rd112, %rd109, 16, 16;
bfi.b64 %rd154, %rd113, %rd106, 32, 32;
setp.lt.u64 %p2, %rd7, %rd154;
@%p2 bra $L__BB7_13;
setp.gt.u64 %p3, %rd7, %rd154;
@%p3 bra $L__BB7_12;
ld.u64 %rd9, [%rd28+32];
setp.lt.u64 %p4, %rd9, %rd155;
@%p4 bra $L__BB7_13;
setp.gt.u64 %p5, %rd9, %rd155;
@%p5 bra $L__BB7_12;
ld.u64 %rd10, [%rd28+24];
setp.lt.u64 %p6, %rd10, %rd156;
@%p6 bra $L__BB7_13;
setp.gt.u64 %p7, %rd10, %rd156;
@%p7 bra $L__BB7_12;
ld.u64 %rd11, [%rd28+16];
setp.lt.u64 %p8, %rd11, %rd157;
@%p8 bra $L__BB7_13;
setp.gt.u64 %p9, %rd11, %rd157;
@%p9 bra $L__BB7_12;
ld.u64 %rd12, [%rd28+8];
setp.lt.u64 %p10, %rd12, %rd158;
@%p10 bra $L__BB7_13;
setp.gt.u64 %p11, %rd12, %rd158;
@%p11 bra $L__BB7_12;
ld.u64 %rd114, [%rd28];
setp.le.u64 %p12, %rd114, %rd159;
@%p12 bra $L__BB7_13;
$L__BB7_12:
// begin inline asm
add.cc.u64 %rd159, %rd159, %rd159;
addc.cc.u64 %rd158, %rd158, %rd158;
addc.cc.u64 %rd157, %rd157, %rd157;
addc.cc.u64 %rd156, %rd156, %rd156;
addc.cc.u64 %rd155, %rd155, %rd155;
addc.u64 %rd154, %rd154, %rd154;
// end inline asm
$L__BB7_13:
ld.u64 %rd145, [%rd28];
ld.u64 %rd146, [%rd28+8];
ld.u64 %rd147, [%rd28+16];
ld.u64 %rd148, [%rd28+24];
ld.u64 %rd149, [%rd28+32];
// begin inline asm
sub.cc.u64 %rd133, %rd159, %rd145;
subc.cc.u64 %rd134, %rd158, %rd146;
subc.cc.u64 %rd135, %rd157, %rd147;
subc.cc.u64 %rd136, %rd156, %rd148;
subc.cc.u64 %rd137, %rd155, %rd149;
subc.u64 %rd138, %rd154, %rd7;
// end inline asm
st.u64 [%rd27], %rd133;
st.u64 [%rd27+8], %rd134;
st.u64 [%rd27+16], %rd135;
st.u64 [%rd27+24], %rd136;
st.u64 [%rd27+32], %rd137;
st.u64 [%rd27+40], %rd138;
bra.uni $L__BB7_16;
$L__BB7_14:
mov.u64 %rd160, 0;
mov.u32 %r4, 0;
$L__BB7_15:
add.s64 %rd152, %rd28, %rd160;
ld.u8 %rs2, [%rd152];
add.s64 %rd153, %rd27, %rd160;
st.u8 [%rd153], %rs2;
add.s64 %rd160, %rd160, 1;
add.s32 %r4, %r4, 1;
setp.lt.u32 %p13, %r4, 48;
@%p13 bra $L__BB7_15;
$L__BB7_16:
ret;
}