//
// 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
// .globl msm6_pixel
.extern .func _Z35blst_p1_add_affines_into_projectiveP7blst_p1PK14blst_p1_affineS3_
(
.param .b64 _Z35blst_p1_add_affines_into_projectiveP7blst_p1PK14blst_p1_affineS3__param_0,
.param .b64 _Z35blst_p1_add_affines_into_projectiveP7blst_p1PK14blst_p1_affineS3__param_1,
.param .b64 _Z35blst_p1_add_affines_into_projectiveP7blst_p1PK14blst_p1_affineS3__param_2
)
;
.extern .func _Z36blst_p1_add_projective_to_projectiveP7blst_p1PKS_S2_
(
.param .b64 _Z36blst_p1_add_projective_to_projectiveP7blst_p1PKS_S2__param_0,
.param .b64 _Z36blst_p1_add_projective_to_projectiveP7blst_p1PKS_S2__param_1,
.param .b64 _Z36blst_p1_add_projective_to_projectiveP7blst_p1PKS_S2__param_2
)
;
.extern .func _Z32blst_p1_add_affine_to_projectiveP7blst_p1PKS_PK14blst_p1_affine
(
.param .b64 _Z32blst_p1_add_affine_to_projectiveP7blst_p1PKS_PK14blst_p1_affine_param_0,
.param .b64 _Z32blst_p1_add_affine_to_projectiveP7blst_p1PKS_PK14blst_p1_affine_param_1,
.param .b64 _Z32blst_p1_add_affine_to_projectiveP7blst_p1PKS_PK14blst_p1_affine_param_2
)
;
.global .align 8 .b8 __nv_static_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_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_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_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_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_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_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_BLS12_377_p0 = -8860621160618917889;
.extern .global .align 8 .b8 BLS12_377_ZERO_PROJECTIVE[144];
.visible .entry msm6_pixel(
.param .u64 msm6_pixel_param_0,
.param .u64 msm6_pixel_param_1,
.param .u64 msm6_pixel_param_2,
.param .u64 msm6_pixel_param_3,
.param .u32 msm6_pixel_param_4
)
{
.local .align 8 .b8 __local_depot0[800];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<13>;
.reg .b16 %rs<3>;
.reg .b32 %r<40>;
.reg .b64 %rd<67>;
.loc 5 8 0
$L__func_begin0:
.loc 5 8 0
mov.u64 %SPL, __local_depot0;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd10, [msm6_pixel_param_0];
ld.param.u64 %rd11, [msm6_pixel_param_1];
ld.param.u64 %rd12, [msm6_pixel_param_2];
ld.param.u64 %rd13, [msm6_pixel_param_3];
ld.param.u32 %r14, [msm6_pixel_param_4];
$L__tmp0:
.loc 5 9 18
mov.u32 %r15, %tid.x;
div.u32 %r16, %r15, 64;
cvt.u64.u32 %rd1, %r16;
$L__tmp1:
.loc 5 10 18
mov.u32 %r17, %tid.x;
cvt.u64.u32 %rd15, %r17;
mul.lo.s64 %rd16, %rd1, 64;
sub.s64 %rd17, %rd15, %rd16;
$L__tmp2:
.loc 5 11 17
cvt.u32.u64 %r18, %rd17;
mov.u64 %rd18, 1;
shl.b64 %rd2, %rd18, %r18;
$L__tmp3:
.loc 5 0 17
mov.u64 %rd14, 0;
add.u64 %rd3, %SP, 0;
.loc 5 14 5
mov.pred %p1, 0;
mov.u64 %rd65, %rd14;
@%p1 bra $L__BB0_2;
bra.uni $L__BB0_1;
$L__BB0_1:
.loc 5 0 5
mov.u64 %rd4, %rd65;
mov.u64 %rd19, BLS12_377_ZERO_PROJECTIVE;
cvta.global.u64 %rd20, %rd19;
add.s64 %rd21, %rd20, %rd4;
ld.u8 %rs1, [%rd21];
add.s64 %rd22, %rd3, %rd4;
st.u8 [%rd22], %rs1;
add.s64 %rd5, %rd4, 1;
setp.lt.u64 %p2, %rd5, 144;
mov.u64 %rd65, %rd5;
@%p2 bra $L__BB0_1;
bra.uni $L__BB0_2;
$L__BB0_2:
.loc 5 16 27
mov.u32 %r19, %ctaid.x;
mul.lo.s32 %r20, %r19, 128;
$L__tmp4:
.loc 5 17 25
mov.u32 %r21, %ctaid.x;
cvt.u64.u32 %rd23, %r21;
shl.b64 %rd24, %rd23, 2;
add.s64 %rd25, %rd13, %rd24;
ld.u32 %r22, [%rd25];
add.s32 %r1, %r20, %r22;
$L__tmp5:
.loc 5 20 35
mov.u32 %r23, 0;
mov.b32 %r2, %r23;
$L__tmp6:
.loc 5 23 21
mov.b32 %r3, %r20;
$L__tmp7:
.loc 5 23 5
mov.u32 %r35, %r2;
$L__tmp8:
mov.u32 %r36, %r3;
$L__tmp9:
bra.uni $L__BB0_3;
$L__BB0_3:
mov.u32 %r5, %r36;
mov.u32 %r4, %r35;
$L__tmp10:
setp.lt.u32 %p3, %r5, %r1;
not.pred %p4, %p3;
@%p4 bra $L__BB0_9;
bra.uni $L__BB0_4;
$L__BB0_4:
$L__tmp11:
.loc 5 24 20
cvt.u64.u32 %rd54, %r5;
shl.b64 %rd55, %rd54, 5;
add.s64 %rd56, %rd12, %rd55;
shl.b64 %rd57, %rd1, 3;
add.s64 %rd58, %rd56, %rd57;
ld.u64 %rd59, [%rd58];
and.b64 %rd60, %rd59, %rd2;
$L__tmp12:
.loc 5 25 9
setp.eq.s64 %p11, %rd60, 0;
not.pred %p12, %p11;
@%p12 bra $L__BB0_6;
bra.uni $L__BB0_5;
$L__BB0_5:
$L__tmp13:
.loc 5 26 13
mov.u32 %r37, %r4;
$L__tmp14:
bra.uni $L__BB0_7;
$L__tmp15:
$L__BB0_6:
.loc 5 28 9
add.s32 %r6, %r4, 1;
$L__tmp16:
cvt.u64.u32 %rd61, %r4;
shl.b64 %rd62, %rd61, 2;
add.u64 %rd63, %SP, 144;
add.s64 %rd64, %rd63, %rd62;
st.u32 [%rd64], %r5;
mov.u32 %r37, %r6;
$L__tmp17:
bra.uni $L__BB0_7;
$L__tmp18:
$L__BB0_7:
.loc 5 23 53
mov.u32 %r7, %r37;
$L__tmp19:
bra.uni $L__BB0_8;
$L__BB0_8:
add.s32 %r8, %r5, 1;
$L__tmp20:
mov.u32 %r35, %r7;
$L__tmp21:
mov.u32 %r36, %r8;
$L__tmp22:
bra.uni $L__BB0_3;
$L__tmp23:
$L__BB0_9:
.loc 5 30 16
mov.u32 %r24, 0;
mov.b32 %r9, %r24;
$L__tmp24:
.loc 5 31 5
mov.u32 %r38, %r9;
$L__tmp25:
bra.uni $L__BB0_10;
$L__BB0_10:
mov.u32 %r10, %r38;
$L__tmp26:
div.u32 %r25, %r4, 2;
mul.lo.s32 %r26, %r25, 2;
setp.lt.u32 %p5, %r10, %r26;
not.pred %p6, %p5;
@%p6 bra $L__BB0_13;
bra.uni $L__BB0_11;
$L__BB0_11:
$L__tmp27:
.loc 5 33 9
cvt.u64.u32 %rd39, %r10;
shl.b64 %rd40, %rd39, 2;
add.u64 %rd41, %SP, 144;
add.s64 %rd42, %rd41, %rd40;
ld.u32 %r32, [%rd42];
cvt.u64.u32 %rd43, %r32;
mul.lo.s64 %rd44, %rd43, 96;
add.s64 %rd45, %rd11, %rd44;
add.s32 %r33, %r10, 1;
cvt.u64.u32 %rd46, %r33;
shl.b64 %rd47, %rd46, 2;
add.s64 %rd48, %rd41, %rd47;
ld.u32 %r34, [%rd48];
cvt.u64.u32 %rd49, %r34;
mul.lo.s64 %rd50, %rd49, 96;
add.s64 %rd51, %rd11, %rd50;
add.u64 %rd52, %SP, 656;
{ // callseq 1, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd52;
.param .b64 param1;
st.param.b64 [param1+0], %rd45;
.param .b64 param2;
st.param.b64 [param2+0], %rd51;
call.uni
_Z35blst_p1_add_affines_into_projectiveP7blst_p1PK14blst_p1_affineS3_,
(
param0,
param1,
param2
);
} // callseq 1
add.u64 %rd53, %SP, 0;
.loc 5 34 9
{ // callseq 2, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd53;
.param .b64 param1;
st.param.b64 [param1+0], %rd53;
.param .b64 param2;
st.param.b64 [param2+0], %rd52;
call.uni
_Z36blst_p1_add_projective_to_projectiveP7blst_p1PKS_S2_,
(
param0,
param1,
param2
);
} // callseq 2
$L__tmp28:
.loc 5 31 48
bra.uni $L__BB0_12;
$L__BB0_12:
add.s32 %r11, %r10, 2;
$L__tmp29:
mov.u32 %r38, %r11;
$L__tmp30:
bra.uni $L__BB0_10;
$L__tmp31:
$L__BB0_13:
.loc 5 36 5
mov.u32 %r39, %r10;
$L__tmp32:
bra.uni $L__BB0_14;
$L__BB0_14:
mov.u32 %r12, %r39;
$L__tmp33:
setp.lt.u32 %p7, %r12, %r4;
not.pred %p8, %p7;
@%p8 bra $L__BB0_17;
bra.uni $L__BB0_15;
$L__BB0_15:
$L__tmp34:
.loc 5 37 9
cvt.u64.u32 %rd31, %r12;
shl.b64 %rd32, %rd31, 2;
add.u64 %rd33, %SP, 144;
add.s64 %rd34, %rd33, %rd32;
ld.u32 %r31, [%rd34];
cvt.u64.u32 %rd35, %r31;
mul.lo.s64 %rd36, %rd35, 96;
add.s64 %rd37, %rd11, %rd36;
add.u64 %rd38, %SP, 0;
{ // callseq 0, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd38;
.param .b64 param1;
st.param.b64 [param1+0], %rd38;
.param .b64 param2;
st.param.b64 [param2+0], %rd37;
call.uni
_Z32blst_p1_add_affine_to_projectiveP7blst_p1PKS_PK14blst_p1_affine,
(
param0,
param1,
param2
);
} // callseq 0
$L__tmp35:
.loc 5 36 38
bra.uni $L__BB0_16;
$L__BB0_16:
add.s32 %r13, %r12, 1;
$L__tmp36:
mov.u32 %r39, %r13;
$L__tmp37:
bra.uni $L__BB0_14;
$L__tmp38:
$L__BB0_17:
.loc 5 40 5
mov.u32 %r27, %tid.x;
mul.lo.s32 %r28, %r27, %r14;
mov.u32 %r29, %ctaid.x;
add.s32 %r30, %r28, %r29;
cvt.u64.u32 %rd27, %r30;
mul.lo.s64 %rd28, %rd27, 144;
add.s64 %rd6, %rd10, %rd28;
mov.u64 %rd26, 0;
add.u64 %rd7, %SP, 0;
mov.pred %p9, 0;
mov.u64 %rd66, %rd26;
@%p9 bra $L__BB0_19;
bra.uni $L__BB0_18;
$L__BB0_18:
.loc 5 0 5
mov.u64 %rd8, %rd66;
add.s64 %rd29, %rd7, %rd8;
ld.u8 %rs2, [%rd29];
add.s64 %rd30, %rd6, %rd8;
st.u8 [%rd30], %rs2;
add.s64 %rd9, %rd8, 1;
setp.lt.u64 %p10, %rd9, 144;
mov.u64 %rd66, %rd9;
@%p10 bra $L__BB0_18;
bra.uni $L__BB0_19;
$L__BB0_19:
.loc 5 41 1
ret;
$L__tmp39:
$L__func_end0:
}
// .globl msm6_collapse_rows
.visible .entry msm6_collapse_rows(
.param .u64 msm6_collapse_rows_param_0,
.param .u64 msm6_collapse_rows_param_1,
.param .u32 msm6_collapse_rows_param_2
)
{
.local .align 8 .b8 __local_depot1[144];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<7>;
.reg .b16 %rs<3>;
.reg .b32 %r<10>;
.reg .b64 %rd<27>;
.loc 5 43 0
$L__func_begin1:
.loc 5 43 0
mov.u64 %SPL, __local_depot1;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd9, [msm6_collapse_rows_param_0];
ld.param.u64 %rd10, [msm6_collapse_rows_param_1];
ld.param.u32 %r6, [msm6_collapse_rows_param_2];
$L__tmp40:
.loc 5 45 19
mov.u32 %r7, %tid.x;
mul.lo.s32 %r1, %r7, %r6;
$L__tmp41:
.loc 5 46 19
add.s32 %r2, %r1, %r6;
$L__tmp42:
.loc 5 0 19
add.u64 %rd1, %SP, 0;
.loc 5 47 5
cvt.u64.u32 %rd12, %r1;
mul.lo.s64 %rd13, %rd12, 144;
add.s64 %rd2, %rd10, %rd13;
mov.u64 %rd11, 0;
mov.pred %p1, 0;
mov.u64 %rd25, %rd11;
@%p1 bra $L__BB1_2;
bra.uni $L__BB1_1;
$L__BB1_1:
.loc 5 0 5
mov.u64 %rd3, %rd25;
add.s64 %rd14, %rd2, %rd3;
ld.u8 %rs1, [%rd14];
add.s64 %rd15, %rd1, %rd3;
st.u8 [%rd15], %rs1;
add.s64 %rd4, %rd3, 1;
setp.lt.u64 %p2, %rd4, 144;
mov.u64 %rd25, %rd4;
@%p2 bra $L__BB1_1;
bra.uni $L__BB1_2;
$L__BB1_2:
$L__tmp43:
.loc 5 49 21
add.s32 %r3, %r1, 1;
$L__tmp44:
.loc 5 49 5
mov.u32 %r9, %r3;
$L__tmp45:
bra.uni $L__BB1_3;
$L__BB1_3:
mov.u32 %r4, %r9;
$L__tmp46:
setp.lt.u32 %p3, %r4, %r2;
not.pred %p4, %p3;
@%p4 bra $L__BB1_6;
bra.uni $L__BB1_4;
$L__BB1_4:
$L__tmp47:
.loc 5 50 9
cvt.u64.u32 %rd21, %r4;
mul.lo.s64 %rd22, %rd21, 144;
add.s64 %rd23, %rd10, %rd22;
add.u64 %rd24, %SP, 0;
{ // callseq 3, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd24;
.param .b64 param1;
st.param.b64 [param1+0], %rd24;
.param .b64 param2;
st.param.b64 [param2+0], %rd23;
call.uni
_Z36blst_p1_add_projective_to_projectiveP7blst_p1PKS_S2_,
(
param0,
param1,
param2
);
} // callseq 3
$L__tmp48:
.loc 5 49 43
bra.uni $L__BB1_5;
$L__BB1_5:
add.s32 %r5, %r4, 1;
$L__tmp49:
mov.u32 %r9, %r5;
$L__tmp50:
bra.uni $L__BB1_3;
$L__tmp51:
$L__BB1_6:
.loc 5 53 5
mov.u32 %r8, %tid.x;
cvt.u64.u32 %rd17, %r8;
mul.lo.s64 %rd18, %rd17, 144;
add.s64 %rd5, %rd9, %rd18;
mov.u64 %rd16, 0;
add.u64 %rd6, %SP, 0;
mov.pred %p5, 0;
mov.u64 %rd26, %rd16;
@%p5 bra $L__BB1_8;
bra.uni $L__BB1_7;
$L__BB1_7:
.loc 5 0 5
mov.u64 %rd7, %rd26;
add.s64 %rd19, %rd6, %rd7;
ld.u8 %rs2, [%rd19];
add.s64 %rd20, %rd5, %rd7;
st.u8 [%rd20], %rs2;
add.s64 %rd8, %rd7, 1;
setp.lt.u64 %p6, %rd8, 144;
mov.u64 %rd26, %rd8;
@%p6 bra $L__BB1_7;
bra.uni $L__BB1_8;
$L__BB1_8:
.loc 5 54 1
ret;
$L__tmp52:
$L__func_end1:
}
.file 1 "/home/p/aleo/snarkVM/algorithms/src/msm/variable_base/blst_377_cuda/./types.h"
.file 2 "/usr/include/x86_64-linux-gnu/bits/types.h"
.file 3 "/usr/include/x86_64-linux-gnu/bits/stdint-uintn.h"
.file 4 "/usr/lib/gcc/x86_64-linux-gnu/10/include/stddef.h"
.file 5 "/home/p/aleo/snarkVM/algorithms/src/msm/variable_base/blst_377_cuda/./msm.cu"
.section .debug_loc
{
.b64 $L__tmp6
.b64 $L__tmp8
.b8 5
.b8 0
.b8 144
.b8 178
.b8 228
.b8 149
.b8 1
.b64 $L__tmp8
.b64 $L__tmp10
.b8 6
.b8 0
.b8 144
.b8 181
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp10
.b64 $L__tmp14
.b8 5
.b8 0
.b8 144
.b8 180
.b8 228
.b8 149
.b8 1
.b64 $L__tmp14
.b64 $L__tmp16
.b8 6
.b8 0
.b8 144
.b8 183
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp16
.b64 $L__tmp17
.b8 5
.b8 0
.b8 144
.b8 182
.b8 228
.b8 149
.b8 1
.b64 $L__tmp17
.b64 $L__tmp19
.b8 6
.b8 0
.b8 144
.b8 183
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp19
.b64 $L__tmp21
.b8 5
.b8 0
.b8 144
.b8 183
.b8 228
.b8 149
.b8 1
.b64 $L__tmp21
.b64 $L__func_end0
.b8 6
.b8 0
.b8 144
.b8 181
.b8 230
.b8 200
.b8 171
.b8 2
.b64 0
.b64 0
.b64 $L__tmp7
.b64 $L__tmp9
.b8 5
.b8 0
.b8 144
.b8 179
.b8 228
.b8 149
.b8 1
.b64 $L__tmp9
.b64 $L__tmp10
.b8 6
.b8 0
.b8 144
.b8 182
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp10
.b64 $L__tmp20
.b8 5
.b8 0
.b8 144
.b8 181
.b8 228
.b8 149
.b8 1
.b64 $L__tmp20
.b64 $L__tmp22
.b8 5
.b8 0
.b8 144
.b8 184
.b8 228
.b8 149
.b8 1
.b64 $L__tmp22
.b64 $L__func_end0
.b8 6
.b8 0
.b8 144
.b8 182
.b8 230
.b8 200
.b8 171
.b8 2
.b64 0
.b64 0
.b64 $L__tmp24
.b64 $L__tmp25
.b8 5
.b8 0
.b8 144
.b8 185
.b8 228
.b8 149
.b8 1
.b64 $L__tmp25
.b64 $L__tmp26
.b8 6
.b8 0
.b8 144
.b8 184
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp26
.b64 $L__tmp29
.b8 6
.b8 0
.b8 144
.b8 176
.b8 226
.b8 200
.b8 171
.b8 2
.b64 $L__tmp29
.b64 $L__tmp30
.b8 6
.b8 0
.b8 144
.b8 177
.b8 226
.b8 200
.b8 171
.b8 2
.b64 $L__tmp30
.b64 $L__tmp32
.b8 6
.b8 0
.b8 144
.b8 184
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp32
.b64 $L__tmp33
.b8 6
.b8 0
.b8 144
.b8 185
.b8 230
.b8 200
.b8 171
.b8 2
.b64 $L__tmp33
.b64 $L__tmp36
.b8 6
.b8 0
.b8 144
.b8 178
.b8 226
.b8 200
.b8 171
.b8 2
.b64 $L__tmp36
.b64 $L__tmp37
.b8 6
.b8 0
.b8 144
.b8 179
.b8 226
.b8 200
.b8 171
.b8 2
.b64 $L__tmp37
.b64 $L__func_end0
.b8 6
.b8 0
.b8 144
.b8 185
.b8 230
.b8 200
.b8 171
.b8 2
.b64 0
.b64 0
.b64 $L__tmp44
.b64 $L__tmp45
.b8 5
.b8 0
.b8 144
.b8 179
.b8 228
.b8 149
.b8 1
.b64 $L__tmp45
.b64 $L__tmp46
.b8 5
.b8 0
.b8 144
.b8 185
.b8 228
.b8 149
.b8 1
.b64 $L__tmp46
.b64 $L__tmp49
.b8 5
.b8 0
.b8 144
.b8 180
.b8 228
.b8 149
.b8 1
.b64 $L__tmp49
.b64 $L__tmp50
.b8 5
.b8 0
.b8 144
.b8 181
.b8 228
.b8 149
.b8 1
.b64 $L__tmp50
.b64 $L__func_end1
.b8 5
.b8 0
.b8 144
.b8 185
.b8 228
.b8 149
.b8 1
.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 52
.b8 0
.b8 3
.b8 8
.b8 73
.b8 19
.b8 63
.b8 12
.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 10
.b8 19
.b8 1
.b8 3
.b8 8
.b8 11
.b8 11
.b8 58
.b8 11
.b8 59
.b8 11
.b8 0
.b8 0
.b8 11
.b8 13
.b8 0
.b8 3
.b8 8
.b8 73
.b8 19
.b8 58
.b8 11
.b8 59
.b8 11
.b8 56
.b8 10
.b8 0
.b8 0
.b8 12
.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 63
.b8 12
.b8 0
.b8 0
.b8 13
.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 14
.b8 11
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 0
.b8 0
.b8 15
.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 16
.b8 52
.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 17
.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 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 1867
.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,109,115,109,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 228
.b8 1
.b8 46
.b8 5
.b8 9
.b8 3
.b64 __nv_static_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_BLS12_377_P
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,49,95,95,51,56,95,116,109,112,120,102,116,95,48,48,50,56,99,53,101,50,95,48,48,48,48,48
.b8 48,48,48,95,55,95,109,115,109,95,99,112,112,49,95,105,105,95,99,97,97,56,53,52,98,57,95,66,76,83,49,50,95,51,55,55,95,80
.b8 0
.b8 3
.b32 240
.b8 4
.b32 281
.b8 6
.b8 0
.b8 5
.b32 245
.b8 6
.b32 259
.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 228
.b8 1
.b8 51
.b8 5
.b8 9
.b8 3
.b64 __nv_static_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_BLS12_377_ONE
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,49,95,95,51,56,95,116,109,112,120,102,116,95,48,48,50,56,99,53,101,50,95,48,48,48,48,48
.b8 48,48,48,95,55,95,109,115,109,95,99,112,112,49,95,105,105,95,99,97,97,56,53,52,98,57,95,66,76,83,49,50,95,51,55,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 228
.b8 1
.b8 52
.b8 5
.b8 9
.b8 3
.b64 __nv_static_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_BLS12_377_R2
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,49,95,95,51,56,95,116,109,112,120,102,116,95,48,48,50,56,99,53,101,50,95,48,48,48,48,48
.b8 48,48,48,95,55,95,109,115,109,95,99,112,112,49,95,105,105,95,99,97,97,56,53,52,98,57,95,66,76,83,49,50,95,51,55,55,95,82,50
.b8 0
.b8 2
.b8 66,76,83,49,50,95,51,55,55,95,112,48
.b8 0
.b32 240
.b8 1
.b8 60
.b8 5
.b8 9
.b8 3
.b64 __nv_static_51__38_tmpxft_0028c5e2_00000000_7_msm_cpp1_ii_caa854b9_BLS12_377_p0
.b8 95,95,110,118,95,115,116,97,116,105,99,95,53,49,95,95,51,56,95,116,109,112,120,102,116,95,48,48,50,56,99,53,101,50,95,48,48,48,48,48
.b8 48,48,48,95,55,95,109,115,109,95,99,112,112,49,95,105,105,95,99,97,97,56,53,52,98,57,95,66,76,83,49,50,95,51,55,55,95,112,48
.b8 0
.b8 9
.b8 66,76,83,49,50,95,51,55,55,95,90,69,82,79,95,80,82,79,74,69,67,84,73,86,69
.b8 0
.b32 710
.b8 1
.b8 1
.b8 61
.b8 5
.b8 9
.b8 3
.b64 BLS12_377_ZERO_PROJECTIVE
.b8 66,76,83,49,50,95,51,55,55,95,90,69,82,79,95,80,82,79,74,69,67,84,73,86,69
.b8 0
.b8 5
.b32 715
.b8 6
.b32 730
.b8 98,108,115,116,95,112,49
.b8 0
.b8 1
.b8 36
.b8 10
.b8 98,108,115,116,95,112,49
.b8 0
.b8 144
.b8 1
.b8 36
.b8 11
.b8 88
.b8 0
.b32 779
.b8 1
.b8 36
.b8 2
.b8 35
.b8 0
.b8 11
.b8 89
.b8 0
.b32 779
.b8 1
.b8 36
.b8 2
.b8 35
.b8 48
.b8 11
.b8 90
.b8 0
.b32 779
.b8 1
.b8 36
.b8 2
.b8 35
.b8 96
.b8 0
.b8 6
.b32 794
.b8 98,108,115,116,95,102,112
.b8 0
.b8 1
.b8 32
.b8 3
.b32 245
.b8 4
.b32 281
.b8 6
.b8 0
.b8 10
.b8 98,108,115,116,95,112,49,95,97,102,102,105,110,101
.b8 0
.b8 96
.b8 1
.b8 35
.b8 11
.b8 88
.b8 0
.b32 779
.b8 1
.b8 35
.b8 2
.b8 35
.b8 0
.b8 11
.b8 89
.b8 0
.b32 779
.b8 1
.b8 35
.b8 2
.b8 35
.b8 48
.b8 0
.b8 6
.b32 806
.b8 98,108,115,116,95,112,49,95,97,102,102,105,110,101
.b8 0
.b8 1
.b8 35
.b8 6
.b32 890
.b8 95,95,117,105,110,116,51,50,95,116
.b8 0
.b8 2
.b8 42
.b8 7
.b8 117,110,115,105,103,110,101,100,32,105,110,116
.b8 0
.b8 7
.b8 4
.b8 6
.b32 872
.b8 117,105,110,116,51,50,95,116
.b8 0
.b8 3
.b8 26
.b8 6
.b32 936
.b8 115,105,122,101,95,116
.b8 0
.b8 4
.b8 209
.b8 7
.b8 117,110,115,105,103,110,101,100,32,108,111,110,103
.b8 0
.b8 7
.b8 8
.b8 12
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b8 109,115,109,54,95,112,105,120,101,108
.b8 0
.b8 109,115,109,54,95,112,105,120,101,108
.b8 0
.b8 5
.b8 8
.b32 1785
.b8 1
.b8 13
.b8 9
.b8 3
.b64 msm6_pixel_param_0
.b8 7
.b8 98,117,99,107,101,116,95,108,105,115,116,115
.b8 0
.b8 5
.b8 8
.b32 1803
.b8 13
.b8 9
.b8 3
.b64 msm6_pixel_param_1
.b8 7
.b8 98,97,115,101,115,95,105,110
.b8 0
.b8 5
.b8 8
.b32 1812
.b8 13
.b8 9
.b8 3
.b64 msm6_pixel_param_2
.b8 7
.b8 115,99,97,108,97,114,115
.b8 0
.b8 5
.b8 8
.b32 1826
.b8 13
.b8 9
.b8 3
.b64 msm6_pixel_param_3
.b8 7
.b8 119,105,110,100,111,119,95,108,101,110,103,116,104,115
.b8 0
.b8 5
.b8 8
.b32 1847
.b8 13
.b8 9
.b8 3
.b64 msm6_pixel_param_4
.b8 7
.b8 119,105,110,100,111,119,95,99,111,117,110,116
.b8 0
.b8 5
.b8 8
.b32 1856
.b8 14
.b64 $L__tmp0
.b64 $L__tmp39
.b8 15
.b8 6
.b8 11
.b8 3
.b64 __local_depot0
.b8 35
.b8 0
.b8 98,117,99,107,101,116
.b8 0
.b8 5
.b8 13
.b32 715
.b8 15
.b8 6
.b8 12
.b8 3
.b64 __local_depot0
.b8 35
.b8 144,1
.b8 97,99,116,105,118,97,116,101,100,95,98,97,115,101,115
.b8 0
.b8 5
.b8 19
.b32 1791
.b8 16
.b8 6
.b8 144
.b8 177
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 105,110,100,101,120
.b8 0
.b8 5
.b8 9
.b32 245
.b8 16
.b8 7
.b8 144
.b8 183
.b8 226
.b8 144
.b8 147
.b8 215
.b8 4
.b8 2
.b8 115,104,105,102,116
.b8 0
.b8 5
.b8 10
.b32 922
.b8 16
.b8 6
.b8 144
.b8 178
.b8 200
.b8 201
.b8 171
.b8 2
.b8 2
.b8 109,97,115,107
.b8 0
.b8 5
.b8 11
.b32 245
.b8 16
.b8 6
.b8 144
.b8 176
.b8 228
.b8 200
.b8 171
.b8 2
.b8 2
.b8 119,105,110,100,111,119,95,115,116,97,114,116
.b8 0
.b8 5
.b8 16
.b32 906
.b8 16
.b8 5
.b8 144
.b8 177
.b8 228
.b8 149
.b8 1
.b8 2
.b8 119,105,110,100,111,119,95,101,110,100
.b8 0
.b8 5
.b8 17
.b32 906
.b8 17
.b32 .debug_loc
.b8 97,99,116,105,118,97,116,101,100,95,98,97,115,101,95,105,110,100,101,120
.b8 0
.b8 5
.b8 20
.b32 906
.b8 17
.b32 .debug_loc+337
.b8 105
.b8 0
.b8 5
.b8 30
.b32 906
.b8 14
.b64 $L__tmp6
.b64 $L__tmp23
.b8 17
.b32 .debug_loc+204
.b8 105
.b8 0
.b8 5
.b8 23
.b32 906
.b8 14
.b64 $L__tmp11
.b64 $L__tmp18
.b8 16
.b8 7
.b8 144
.b8 176
.b8 236
.b8 144
.b8 147
.b8 215
.b8 4
.b8 2
.b8 98,105,116
.b8 0
.b8 5
.b8 24
.b32 245
.b8 0
.b8 0
.b8 14
.b64 $L__tmp27
.b64 $L__tmp28
.b8 15
.b8 6
.b8 12
.b8 3
.b64 __local_depot0
.b8 35
.b8 144,5
.b8 105,110,116,101,114,109,101,100,105,97,116,101
.b8 0
.b8 5
.b8 32
.b32 715
.b8 0
.b8 0
.b8 0
.b8 12
.b64 $L__func_begin1
.b64 $L__func_end1
.b8 1
.b8 156
.b8 109,115,109,54,95,99,111,108,108,97,112,115,101,95,114,111,119,115
.b8 0
.b8 109,115,109,54,95,99,111,108,108,97,112,115,101,95,114,111,119,115
.b8 0
.b8 5
.b8 43
.b32 1785
.b8 1
.b8 13
.b8 9
.b8 3
.b64 msm6_collapse_rows_param_0
.b8 7
.b8 116,97,114,103,101,116
.b8 0
.b8 5
.b8 43
.b32 1803
.b8 13
.b8 9
.b8 3
.b64 msm6_collapse_rows_param_1
.b8 7
.b8 98,117,99,107,101,116,95,108,105,115,116,115
.b8 0
.b8 5
.b8 43
.b32 1861
.b8 13
.b8 9
.b8 3
.b64 msm6_collapse_rows_param_2
.b8 7
.b8 119,105,110,100,111,119,95,99,111,117,110,116
.b8 0
.b8 5
.b8 43
.b32 1856
.b8 14
.b64 $L__tmp40
.b64 $L__tmp52
.b8 15
.b8 6
.b8 11
.b8 3
.b64 __local_depot1
.b8 35
.b8 0
.b8 116,101,109,112,95,116,97,114,103,101,116
.b8 0
.b8 5
.b8 44
.b32 715
.b8 16
.b8 5
.b8 144
.b8 177
.b8 228
.b8 149
.b8 1
.b8 2
.b8 98,97,115,101
.b8 0
.b8 5
.b8 45
.b32 906
.b8 16
.b8 5
.b8 144
.b8 178
.b8 228
.b8 149
.b8 1
.b8 2
.b8 116,101,114,109
.b8 0
.b8 5
.b8 46
.b32 906
.b8 14
.b64 $L__tmp43
.b64 $L__tmp51
.b8 17
.b32 .debug_loc+568
.b8 105
.b8 0
.b8 5
.b8 49
.b32 906
.b8 0
.b8 0
.b8 0
.b8 18
.b8 118,111,105,100
.b8 0
.b8 3
.b32 906
.b8 4
.b32 281
.b8 128
.b8 0
.b8 19
.b32 715
.b32 12
.b8 19
.b32 1821
.b32 12
.b8 5
.b32 850
.b8 19
.b32 1835
.b32 12
.b8 3
.b32 240
.b8 4
.b32 281
.b8 4
.b8 0
.b8 19
.b32 1856
.b32 12
.b8 5
.b32 906
.b8 19
.b32 710
.b32 12
.b8 0
}
.section .debug_macinfo
{
.b8 0
}