tyler.durden/naive-hashcat
1
0
Fork 0
mirror of https://github.com/brannondorsey/naive-hashcat.git synced 2025-07-31 01:34:51 +02:00
Code Issues Packages Projects Releases Wiki Activity
naive-hashcat/hashcat-3.6.0/kernels/m00400.7607c47d.kernel
Brannon Dorsey fb5d5b8d78 First commit 🐈
2017-07-17 23:41:44 -05:00

1570 lines
48 KiB
Plaintext
Raw Permalink Blame History

This file contains invisible Unicode characters

This file contains invisible Unicode characters that are indistinguishable to humans but may be processed differently by a computer. If you think that this is intentional, you can safely ignore this warning. Use the Escape button to reveal them.

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-22053397
// Driver 375.66
// Based on LLVM 3.4svn
//
.version 5.0
.target sm_61, texmode_independent
.address_size 64
// .globl gpu_memset
.entry gpu_memset(
.param .u64 .ptr .global .align 16 gpu_memset_param_0,
.param .u32 gpu_memset_param_1,
.param .u32 gpu_memset_param_2
)
{
.reg .pred %p<2>;
.reg .b32 %r<9>;
.reg .b64 %rd<4>;
ld.param.u64 %rd1, [gpu_memset_param_0];
ld.param.u32 %r2, [gpu_memset_param_1];
ld.param.u32 %r3, [gpu_memset_param_2];
mov.b32 %r4, %envreg3;
mov.u32 %r5, %ctaid.x;
mov.u32 %r6, %ntid.x;
mad.lo.s32 %r7, %r5, %r6, %r4;
mov.u32 %r8, %tid.x;
add.s32 %r1, %r7, %r8;
setp.ge.u32 %p1, %r1, %r3;
@%p1 bra BB0_2;
mul.wide.u32 %rd2, %r1, 16;
add.s64 %rd3, %rd1, %rd2;
st.global.v4.u32 [%rd3], {%r2, %r2, %r2, %r2};
BB0_2:
ret;
}
// .globl m00400_init
.entry m00400_init(
.param .u64 .ptr .global .align 4 m00400_init_param_0,
.param .u64 .ptr .global .align 4 m00400_init_param_1,
.param .u64 .ptr .global .align 4 m00400_init_param_2,
.param .u64 .ptr .global .align 4 m00400_init_param_3,
.param .u64 .ptr .global .align 4 m00400_init_param_4,
.param .u64 .ptr .global .align 1 m00400_init_param_5,
.param .u64 .ptr .global .align 4 m00400_init_param_6,
.param .u64 .ptr .global .align 4 m00400_init_param_7,
.param .u64 .ptr .global .align 4 m00400_init_param_8,
.param .u64 .ptr .global .align 4 m00400_init_param_9,
.param .u64 .ptr .global .align 4 m00400_init_param_10,
.param .u64 .ptr .global .align 4 m00400_init_param_11,
.param .u64 .ptr .global .align 4 m00400_init_param_12,
.param .u64 .ptr .global .align 4 m00400_init_param_13,
.param .u64 .ptr .global .align 4 m00400_init_param_14,
.param .u64 .ptr .global .align 4 m00400_init_param_15,
.param .u64 .ptr .global .align 4 m00400_init_param_16,
.param .u64 .ptr .global .align 4 m00400_init_param_17,
.param .u64 .ptr .global .align 1 m00400_init_param_18,
.param .u64 .ptr .global .align 4 m00400_init_param_19,
.param .u64 .ptr .global .align 4 m00400_init_param_20,
.param .u64 .ptr .global .align 4 m00400_init_param_21,
.param .u64 .ptr .global .align 4 m00400_init_param_22,
.param .u64 .ptr .global .align 4 m00400_init_param_23,
.param .u32 m00400_init_param_24,
.param .u32 m00400_init_param_25,
.param .u32 m00400_init_param_26,
.param .u32 m00400_init_param_27,
.param .u32 m00400_init_param_28,
.param .u32 m00400_init_param_29,
.param .u32 m00400_init_param_30,
.param .u32 m00400_init_param_31,
.param .u32 m00400_init_param_32,
.param .u32 m00400_init_param_33,
.param .u32 m00400_init_param_34
)
{
.reg .pred %p<21>;
.reg .b32 %r<542>;
.reg .b64 %rd<10>;
ld.param.u64 %rd1, [m00400_init_param_0];
ld.param.u64 %rd2, [m00400_init_param_4];
ld.param.u64 %rd3, [m00400_init_param_17];
ld.param.u32 %r2, [m00400_init_param_27];
ld.param.u32 %r3, [m00400_init_param_34];
mov.b32 %r4, %envreg3;
mov.u32 %r5, %ctaid.x;
mov.u32 %r6, %ntid.x;
mad.lo.s32 %r7, %r5, %r6, %r4;
mov.u32 %r8, %tid.x;
add.s32 %r1, %r7, %r8;
setp.ge.u32 %p1, %r1, %r3;
@%p1 bra BB1_2;
mul.wide.u32 %rd4, %r1, 80;
add.s64 %rd5, %rd1, %rd4;
mul.wide.u32 %rd6, %r2, 180;
add.s64 %rd7, %rd3, %rd6;
ld.global.u32 %r9, [%rd5+64];
add.s32 %r10, %r9, 8;
shl.b32 %r11, %r10, 3;
and.b32 %r12, %r11, 24;
mov.u32 %r13, 128;
shl.b32 %r14, %r13, %r12;
setp.lt.u32 %p2, %r10, 4;
selp.b32 %r15, %r14, 0, %p2;
ld.global.u32 %r16, [%rd7];
or.b32 %r17, %r15, %r16;
setp.gt.u32 %p3, %r10, 3;
setp.gt.u32 %p4, %r9, -9;
and.pred %p5, %p3, %p4;
selp.b32 %r18, %r14, 0, %p5;
ld.global.u32 %r19, [%rd7+4];
or.b32 %r20, %r18, %r19;
setp.lt.u32 %p6, %r9, -8;
setp.lt.u32 %p7, %r10, 12;
and.pred %p8, %p6, %p7;
selp.b32 %r21, %r14, 0, %p8;
ld.global.u32 %r22, [%rd5];
or.b32 %r23, %r21, %r22;
and.b32 %r24, %r10, -4;
setp.eq.s32 %p9, %r24, 12;
selp.b32 %r25, %r14, 0, %p9;
ld.global.u32 %r26, [%rd5+4];
or.b32 %r27, %r25, %r26;
setp.eq.s32 %p10, %r24, 16;
selp.b32 %r28, %r14, 0, %p10;
ld.global.u32 %r29, [%rd5+8];
or.b32 %r30, %r28, %r29;
setp.eq.s32 %p11, %r24, 20;
selp.b32 %r31, %r14, 0, %p11;
ld.global.u32 %r32, [%rd5+12];
or.b32 %r33, %r31, %r32;
setp.eq.s32 %p12, %r24, 24;
selp.b32 %r34, %r14, 0, %p12;
ld.global.u32 %r35, [%rd5+16];
or.b32 %r36, %r34, %r35;
setp.eq.s32 %p13, %r24, 28;
selp.b32 %r37, %r14, 0, %p13;
ld.global.u32 %r38, [%rd5+20];
or.b32 %r39, %r37, %r38;
setp.eq.s32 %p14, %r24, 32;
selp.b32 %r40, %r14, 0, %p14;
ld.global.u32 %r41, [%rd5+24];
or.b32 %r42, %r40, %r41;
setp.eq.s32 %p15, %r24, 36;
selp.b32 %r43, %r14, 0, %p15;
ld.global.u32 %r44, [%rd5+28];
or.b32 %r45, %r43, %r44;
setp.eq.s32 %p16, %r24, 40;
selp.b32 %r46, %r14, 0, %p16;
ld.global.u32 %r47, [%rd5+32];
or.b32 %r48, %r46, %r47;
setp.eq.s32 %p17, %r24, 44;
selp.b32 %r49, %r14, 0, %p17;
ld.global.u32 %r50, [%rd5+36];
or.b32 %r51, %r49, %r50;
setp.eq.s32 %p18, %r24, 48;
selp.b32 %r52, %r14, 0, %p18;
setp.eq.s32 %p19, %r24, 52;
selp.b32 %r53, %r14, 0, %p19;
setp.eq.s32 %p20, %r24, 56;
selp.b32 %r54, %r14, 0, %p20;
or.b32 %r55, %r54, %r11;
add.s32 %r56, %r17, -680876937;
shf.l.wrap.b32 %r57, %r56, %r56, 7;
add.s32 %r58, %r57, -271733879;
and.b32 %r59, %r58, 2004318071;
xor.b32 %r60, %r59, -1732584194;
add.s32 %r61, %r20, %r60;
add.s32 %r62, %r61, -117830708;
shf.l.wrap.b32 %r63, %r62, %r62, 12;
add.s32 %r64, %r63, %r58;
xor.b32 %r65, %r58, -271733879;
and.b32 %r66, %r64, %r65;
xor.b32 %r67, %r66, -271733879;
add.s32 %r68, %r23, %r67;
add.s32 %r69, %r68, -1126478375;
shf.l.wrap.b32 %r70, %r69, %r69, 17;
add.s32 %r71, %r70, %r64;
xor.b32 %r72, %r64, %r58;
and.b32 %r73, %r71, %r72;
xor.b32 %r74, %r73, %r58;
add.s32 %r75, %r27, %r74;
add.s32 %r76, %r75, -1316259209;
shf.l.wrap.b32 %r77, %r76, %r76, 22;
add.s32 %r78, %r77, %r71;
add.s32 %r79, %r57, %r30;
xor.b32 %r80, %r71, %r64;
and.b32 %r81, %r78, %r80;
xor.b32 %r82, %r81, %r64;
add.s32 %r83, %r79, %r82;
add.s32 %r84, %r83, -448152776;
shf.l.wrap.b32 %r85, %r84, %r84, 7;
add.s32 %r86, %r85, %r78;
add.s32 %r87, %r64, %r33;
xor.b32 %r88, %r78, %r71;
and.b32 %r89, %r86, %r88;
xor.b32 %r90, %r89, %r71;
add.s32 %r91, %r87, %r90;
add.s32 %r92, %r91, 1200080426;
shf.l.wrap.b32 %r93, %r92, %r92, 12;
add.s32 %r94, %r93, %r86;
add.s32 %r95, %r71, %r36;
xor.b32 %r96, %r86, %r78;
and.b32 %r97, %r94, %r96;
xor.b32 %r98, %r97, %r78;
add.s32 %r99, %r95, %r98;
add.s32 %r100, %r99, -1473231341;
shf.l.wrap.b32 %r101, %r100, %r100, 17;
add.s32 %r102, %r101, %r94;
xor.b32 %r103, %r94, %r86;
and.b32 %r104, %r102, %r103;
xor.b32 %r105, %r104, %r86;
add.s32 %r106, %r39, %r78;
add.s32 %r107, %r106, %r105;
add.s32 %r108, %r107, -45705983;
shf.l.wrap.b32 %r109, %r108, %r108, 22;
add.s32 %r110, %r109, %r102;
xor.b32 %r111, %r102, %r94;
and.b32 %r112, %r110, %r111;
xor.b32 %r113, %r112, %r94;
add.s32 %r114, %r42, %r86;
add.s32 %r115, %r114, %r113;
add.s32 %r116, %r115, 1770035416;
shf.l.wrap.b32 %r117, %r116, %r116, 7;
add.s32 %r118, %r117, %r110;
xor.b32 %r119, %r110, %r102;
and.b32 %r120, %r118, %r119;
xor.b32 %r121, %r120, %r102;
add.s32 %r122, %r45, %r94;
add.s32 %r123, %r122, %r121;
add.s32 %r124, %r123, -1958414417;
shf.l.wrap.b32 %r125, %r124, %r124, 12;
add.s32 %r126, %r125, %r118;
xor.b32 %r127, %r118, %r110;
and.b32 %r128, %r126, %r127;
xor.b32 %r129, %r128, %r110;
add.s32 %r130, %r48, %r102;
add.s32 %r131, %r130, %r129;
add.s32 %r132, %r131, -42063;
shf.l.wrap.b32 %r133, %r132, %r132, 17;
add.s32 %r134, %r133, %r126;
xor.b32 %r135, %r126, %r118;
and.b32 %r136, %r134, %r135;
xor.b32 %r137, %r136, %r118;
add.s32 %r138, %r51, %r110;
add.s32 %r139, %r138, %r137;
add.s32 %r140, %r139, -1990404162;
shf.l.wrap.b32 %r141, %r140, %r140, 22;
add.s32 %r142, %r141, %r134;
xor.b32 %r143, %r134, %r126;
and.b32 %r144, %r142, %r143;
xor.b32 %r145, %r144, %r126;
add.s32 %r146, %r52, %r118;
add.s32 %r147, %r146, %r145;
add.s32 %r148, %r147, 1804603682;
shf.l.wrap.b32 %r149, %r148, %r148, 7;
add.s32 %r150, %r149, %r142;
xor.b32 %r151, %r142, %r134;
and.b32 %r152, %r150, %r151;
xor.b32 %r153, %r152, %r134;
add.s32 %r154, %r53, %r126;
add.s32 %r155, %r154, %r153;
add.s32 %r156, %r155, -40341101;
shf.l.wrap.b32 %r157, %r156, %r156, 12;
add.s32 %r158, %r157, %r150;
xor.b32 %r159, %r150, %r142;
and.b32 %r160, %r158, %r159;
xor.b32 %r161, %r160, %r142;
add.s32 %r162, %r55, %r134;
add.s32 %r163, %r162, %r161;
add.s32 %r164, %r163, -1502002290;
shf.l.wrap.b32 %r165, %r164, %r164, 17;
add.s32 %r166, %r165, %r158;
xor.b32 %r167, %r158, %r150;
and.b32 %r168, %r166, %r167;
xor.b32 %r169, %r168, %r150;
add.s32 %r170, %r142, %r169;
add.s32 %r171, %r170, 1236535329;
shf.l.wrap.b32 %r172, %r171, %r171, 22;
add.s32 %r173, %r172, %r166;
xor.b32 %r174, %r173, %r166;
and.b32 %r175, %r174, %r158;
xor.b32 %r176, %r175, %r166;
add.s32 %r177, %r20, %r150;
add.s32 %r178, %r177, %r176;
add.s32 %r179, %r178, -165796510;
shf.l.wrap.b32 %r180, %r179, %r179, 5;
add.s32 %r181, %r180, %r173;
xor.b32 %r182, %r181, %r173;
and.b32 %r183, %r182, %r166;
xor.b32 %r184, %r183, %r173;
add.s32 %r185, %r36, %r158;
add.s32 %r186, %r185, %r184;
add.s32 %r187, %r186, -1069501632;
shf.l.wrap.b32 %r188, %r187, %r187, 9;
add.s32 %r189, %r188, %r181;
xor.b32 %r190, %r189, %r181;
and.b32 %r191, %r190, %r173;
xor.b32 %r192, %r191, %r181;
add.s32 %r193, %r51, %r166;
add.s32 %r194, %r193, %r192;
add.s32 %r195, %r194, 643717713;
shf.l.wrap.b32 %r196, %r195, %r195, 14;
add.s32 %r197, %r196, %r189;
xor.b32 %r198, %r197, %r189;
and.b32 %r199, %r198, %r181;
xor.b32 %r200, %r199, %r189;
add.s32 %r201, %r17, %r173;
add.s32 %r202, %r201, %r200;
add.s32 %r203, %r202, -373897302;
shf.l.wrap.b32 %r204, %r203, %r203, 20;
add.s32 %r205, %r204, %r197;
xor.b32 %r206, %r205, %r197;
and.b32 %r207, %r206, %r189;
xor.b32 %r208, %r207, %r197;
add.s32 %r209, %r33, %r181;
add.s32 %r210, %r209, %r208;
add.s32 %r211, %r210, -701558691;
shf.l.wrap.b32 %r212, %r211, %r211, 5;
add.s32 %r213, %r212, %r205;
xor.b32 %r214, %r213, %r205;
and.b32 %r215, %r214, %r197;
xor.b32 %r216, %r215, %r205;
add.s32 %r217, %r48, %r189;
add.s32 %r218, %r217, %r216;
add.s32 %r219, %r218, 38016083;
shf.l.wrap.b32 %r220, %r219, %r219, 9;
add.s32 %r221, %r220, %r213;
xor.b32 %r222, %r221, %r213;
and.b32 %r223, %r222, %r205;
xor.b32 %r224, %r223, %r213;
add.s32 %r225, %r197, %r224;
add.s32 %r226, %r225, -660478335;
shf.l.wrap.b32 %r227, %r226, %r226, 14;
add.s32 %r228, %r227, %r221;
xor.b32 %r229, %r228, %r221;
and.b32 %r230, %r229, %r213;
xor.b32 %r231, %r230, %r221;
add.s32 %r232, %r30, %r205;
add.s32 %r233, %r232, %r231;
add.s32 %r234, %r233, -405537848;
shf.l.wrap.b32 %r235, %r234, %r234, 20;
add.s32 %r236, %r235, %r228;
xor.b32 %r237, %r236, %r228;
and.b32 %r238, %r237, %r221;
xor.b32 %r239, %r238, %r228;
add.s32 %r240, %r45, %r213;
add.s32 %r241, %r240, %r239;
add.s32 %r242, %r241, 568446438;
shf.l.wrap.b32 %r243, %r242, %r242, 5;
add.s32 %r244, %r243, %r236;
xor.b32 %r245, %r244, %r236;
and.b32 %r246, %r245, %r228;
xor.b32 %r247, %r246, %r236;
add.s32 %r248, %r55, %r221;
add.s32 %r249, %r248, %r247;
add.s32 %r250, %r249, -1019803690;
shf.l.wrap.b32 %r251, %r250, %r250, 9;
add.s32 %r252, %r251, %r244;
xor.b32 %r253, %r252, %r244;
and.b32 %r254, %r253, %r236;
xor.b32 %r255, %r254, %r244;
add.s32 %r256, %r27, %r228;
add.s32 %r257, %r256, %r255;
add.s32 %r258, %r257, -187363961;
shf.l.wrap.b32 %r259, %r258, %r258, 14;
add.s32 %r260, %r259, %r252;
xor.b32 %r261, %r260, %r252;
and.b32 %r262, %r261, %r244;
xor.b32 %r263, %r262, %r252;
add.s32 %r264, %r42, %r236;
add.s32 %r265, %r264, %r263;
add.s32 %r266, %r265, 1163531501;
shf.l.wrap.b32 %r267, %r266, %r266, 20;
add.s32 %r268, %r267, %r260;
xor.b32 %r269, %r268, %r260;
and.b32 %r270, %r269, %r252;
xor.b32 %r271, %r270, %r260;
add.s32 %r272, %r53, %r244;
add.s32 %r273, %r272, %r271;
add.s32 %r274, %r273, -1444681467;
shf.l.wrap.b32 %r275, %r274, %r274, 5;
add.s32 %r276, %r275, %r268;
xor.b32 %r277, %r276, %r268;
and.b32 %r278, %r277, %r260;
xor.b32 %r279, %r278, %r268;
add.s32 %r280, %r23, %r252;
add.s32 %r281, %r280, %r279;
add.s32 %r282, %r281, -51403784;
shf.l.wrap.b32 %r283, %r282, %r282, 9;
add.s32 %r284, %r283, %r276;
xor.b32 %r285, %r284, %r276;
and.b32 %r286, %r285, %r268;
xor.b32 %r287, %r286, %r276;
add.s32 %r288, %r39, %r260;
add.s32 %r289, %r288, %r287;
add.s32 %r290, %r289, 1735328473;
shf.l.wrap.b32 %r291, %r290, %r290, 14;
add.s32 %r292, %r291, %r284;
xor.b32 %r293, %r292, %r284;
and.b32 %r294, %r293, %r276;
xor.b32 %r295, %r294, %r284;
add.s32 %r296, %r52, %r268;
add.s32 %r297, %r296, %r295;
add.s32 %r298, %r297, -1926607734;
shf.l.wrap.b32 %r299, %r298, %r298, 20;
add.s32 %r300, %r299, %r292;
xor.b32 %r301, %r293, %r300;
add.s32 %r302, %r33, %r276;
add.s32 %r303, %r302, %r301;
add.s32 %r304, %r303, -378558;
shf.l.wrap.b32 %r305, %r304, %r304, 4;
add.s32 %r306, %r305, %r300;
xor.b32 %r307, %r300, %r292;
xor.b32 %r308, %r307, %r306;
add.s32 %r309, %r42, %r284;
add.s32 %r310, %r309, %r308;
add.s32 %r311, %r310, -2022574463;
shf.l.wrap.b32 %r312, %r311, %r311, 11;
add.s32 %r313, %r312, %r306;
xor.b32 %r314, %r306, %r300;
xor.b32 %r315, %r314, %r313;
add.s32 %r316, %r51, %r292;
add.s32 %r317, %r316, %r315;
add.s32 %r318, %r317, 1839030562;
shf.l.wrap.b32 %r319, %r318, %r318, 16;
add.s32 %r320, %r319, %r313;
xor.b32 %r321, %r313, %r306;
xor.b32 %r322, %r321, %r320;
add.s32 %r323, %r55, %r300;
add.s32 %r324, %r323, %r322;
add.s32 %r325, %r324, -35309556;
shf.l.wrap.b32 %r326, %r325, %r325, 23;
add.s32 %r327, %r326, %r320;
xor.b32 %r328, %r320, %r313;
xor.b32 %r329, %r328, %r327;
add.s32 %r330, %r20, %r306;
add.s32 %r331, %r330, %r329;
add.s32 %r332, %r331, -1530992060;
shf.l.wrap.b32 %r333, %r332, %r332, 4;
add.s32 %r334, %r333, %r327;
xor.b32 %r335, %r327, %r320;
xor.b32 %r336, %r335, %r334;
add.s32 %r337, %r30, %r313;
add.s32 %r338, %r337, %r336;
add.s32 %r339, %r338, 1272893353;
shf.l.wrap.b32 %r340, %r339, %r339, 11;
add.s32 %r341, %r340, %r334;
xor.b32 %r342, %r334, %r327;
xor.b32 %r343, %r342, %r341;
add.s32 %r344, %r39, %r320;
add.s32 %r345, %r344, %r343;
add.s32 %r346, %r345, -155497632;
shf.l.wrap.b32 %r347, %r346, %r346, 16;
add.s32 %r348, %r347, %r341;
xor.b32 %r349, %r341, %r334;
xor.b32 %r350, %r349, %r348;
add.s32 %r351, %r48, %r327;
add.s32 %r352, %r351, %r350;
add.s32 %r353, %r352, -1094730640;
shf.l.wrap.b32 %r354, %r353, %r353, 23;
add.s32 %r355, %r354, %r348;
xor.b32 %r356, %r348, %r341;
xor.b32 %r357, %r356, %r355;
add.s32 %r358, %r53, %r334;
add.s32 %r359, %r358, %r357;
add.s32 %r360, %r359, 681279174;
shf.l.wrap.b32 %r361, %r360, %r360, 4;
add.s32 %r362, %r361, %r355;
xor.b32 %r363, %r355, %r348;
xor.b32 %r364, %r363, %r362;
add.s32 %r365, %r17, %r341;
add.s32 %r366, %r365, %r364;
add.s32 %r367, %r366, -358537222;
shf.l.wrap.b32 %r368, %r367, %r367, 11;
add.s32 %r369, %r368, %r362;
xor.b32 %r370, %r362, %r355;
xor.b32 %r371, %r370, %r369;
add.s32 %r372, %r27, %r348;
add.s32 %r373, %r372, %r371;
add.s32 %r374, %r373, -722521979;
shf.l.wrap.b32 %r375, %r374, %r374, 16;
add.s32 %r376, %r375, %r369;
xor.b32 %r377, %r369, %r362;
xor.b32 %r378, %r377, %r376;
add.s32 %r379, %r36, %r355;
add.s32 %r380, %r379, %r378;
add.s32 %r381, %r380, 76029189;
shf.l.wrap.b32 %r382, %r381, %r381, 23;
add.s32 %r383, %r382, %r376;
xor.b32 %r384, %r376, %r369;
xor.b32 %r385, %r384, %r383;
add.s32 %r386, %r45, %r362;
add.s32 %r387, %r386, %r385;
add.s32 %r388, %r387, -640364487;
shf.l.wrap.b32 %r389, %r388, %r388, 4;
add.s32 %r390, %r389, %r383;
xor.b32 %r391, %r383, %r376;
xor.b32 %r392, %r391, %r390;
add.s32 %r393, %r52, %r369;
add.s32 %r394, %r393, %r392;
add.s32 %r395, %r394, -421815835;
shf.l.wrap.b32 %r396, %r395, %r395, 11;
add.s32 %r397, %r396, %r390;
xor.b32 %r398, %r390, %r383;
xor.b32 %r399, %r398, %r397;
add.s32 %r400, %r376, %r399;
add.s32 %r401, %r400, 530742520;
shf.l.wrap.b32 %r402, %r401, %r401, 16;
add.s32 %r403, %r402, %r397;
xor.b32 %r404, %r397, %r390;
xor.b32 %r405, %r404, %r403;
add.s32 %r406, %r23, %r383;
add.s32 %r407, %r406, %r405;
add.s32 %r408, %r407, -995338651;
shf.l.wrap.b32 %r409, %r408, %r408, 23;
add.s32 %r410, %r409, %r403;
not.b32 %r411, %r397;
or.b32 %r412, %r410, %r411;
xor.b32 %r413, %r412, %r403;
add.s32 %r414, %r17, %r390;
add.s32 %r415, %r414, %r413;
add.s32 %r416, %r415, -198630844;
shf.l.wrap.b32 %r417, %r416, %r416, 6;
add.s32 %r418, %r417, %r410;
not.b32 %r419, %r403;
or.b32 %r420, %r418, %r419;
xor.b32 %r421, %r420, %r410;
add.s32 %r422, %r39, %r397;
add.s32 %r423, %r422, %r421;
add.s32 %r424, %r423, 1126891415;
shf.l.wrap.b32 %r425, %r424, %r424, 10;
add.s32 %r426, %r425, %r418;
not.b32 %r427, %r410;
or.b32 %r428, %r426, %r427;
xor.b32 %r429, %r428, %r418;
add.s32 %r430, %r55, %r403;
add.s32 %r431, %r430, %r429;
add.s32 %r432, %r431, -1416354905;
shf.l.wrap.b32 %r433, %r432, %r432, 15;
add.s32 %r434, %r433, %r426;
not.b32 %r435, %r418;
or.b32 %r436, %r434, %r435;
xor.b32 %r437, %r436, %r426;
add.s32 %r438, %r33, %r410;
add.s32 %r439, %r438, %r437;
add.s32 %r440, %r439, -57434055;
shf.l.wrap.b32 %r441, %r440, %r440, 21;
add.s32 %r442, %r441, %r434;
not.b32 %r443, %r426;
or.b32 %r444, %r442, %r443;
xor.b32 %r445, %r444, %r434;
add.s32 %r446, %r52, %r418;
add.s32 %r447, %r446, %r445;
add.s32 %r448, %r447, 1700485571;
shf.l.wrap.b32 %r449, %r448, %r448, 6;
add.s32 %r450, %r449, %r442;
not.b32 %r451, %r434;
or.b32 %r452, %r450, %r451;
xor.b32 %r453, %r452, %r442;
add.s32 %r454, %r27, %r426;
add.s32 %r455, %r454, %r453;
add.s32 %r456, %r455, -1894986606;
shf.l.wrap.b32 %r457, %r456, %r456, 10;
add.s32 %r458, %r457, %r450;
not.b32 %r459, %r442;
or.b32 %r460, %r458, %r459;
xor.b32 %r461, %r460, %r450;
add.s32 %r462, %r48, %r434;
add.s32 %r463, %r462, %r461;
add.s32 %r464, %r463, -1051523;
shf.l.wrap.b32 %r465, %r464, %r464, 15;
add.s32 %r466, %r465, %r458;
not.b32 %r467, %r450;
or.b32 %r468, %r466, %r467;
xor.b32 %r469, %r468, %r458;
add.s32 %r470, %r20, %r442;
add.s32 %r471, %r470, %r469;
add.s32 %r472, %r471, -2054922799;
shf.l.wrap.b32 %r473, %r472, %r472, 21;
add.s32 %r474, %r473, %r466;
not.b32 %r475, %r458;
or.b32 %r476, %r474, %r475;
xor.b32 %r477, %r476, %r466;
add.s32 %r478, %r42, %r450;
add.s32 %r479, %r478, %r477;
add.s32 %r480, %r479, 1873313359;
shf.l.wrap.b32 %r481, %r480, %r480, 6;
add.s32 %r482, %r481, %r474;
not.b32 %r483, %r466;
or.b32 %r484, %r482, %r483;
xor.b32 %r485, %r484, %r474;
add.s32 %r486, %r458, %r485;
add.s32 %r487, %r486, -30611744;
shf.l.wrap.b32 %r488, %r487, %r487, 10;
add.s32 %r489, %r488, %r482;
not.b32 %r490, %r474;
or.b32 %r491, %r489, %r490;
xor.b32 %r492, %r491, %r482;
add.s32 %r493, %r36, %r466;
add.s32 %r494, %r493, %r492;
add.s32 %r495, %r494, -1560198380;
shf.l.wrap.b32 %r496, %r495, %r495, 15;
add.s32 %r497, %r496, %r489;
not.b32 %r498, %r482;
or.b32 %r499, %r497, %r498;
xor.b32 %r500, %r499, %r489;
add.s32 %r501, %r53, %r474;
add.s32 %r502, %r501, %r500;
add.s32 %r503, %r502, 1309151649;
shf.l.wrap.b32 %r504, %r503, %r503, 21;
add.s32 %r505, %r504, %r497;
not.b32 %r506, %r489;
or.b32 %r507, %r505, %r506;
xor.b32 %r508, %r507, %r497;
add.s32 %r509, %r30, %r482;
add.s32 %r510, %r509, %r508;
add.s32 %r511, %r510, -145523070;
shf.l.wrap.b32 %r512, %r511, %r511, 6;
add.s32 %r513, %r512, %r505;
not.b32 %r514, %r497;
or.b32 %r515, %r513, %r514;
xor.b32 %r516, %r515, %r505;
add.s32 %r517, %r51, %r489;
add.s32 %r518, %r517, %r516;
add.s32 %r519, %r518, -1120210379;
shf.l.wrap.b32 %r520, %r519, %r519, 10;
add.s32 %r521, %r520, %r513;
not.b32 %r522, %r505;
or.b32 %r523, %r521, %r522;
xor.b32 %r524, %r523, %r513;
add.s32 %r525, %r23, %r497;
add.s32 %r526, %r525, %r524;
add.s32 %r527, %r526, 718787259;
shf.l.wrap.b32 %r528, %r527, %r527, 15;
add.s32 %r529, %r528, %r521;
not.b32 %r530, %r513;
or.b32 %r531, %r529, %r530;
xor.b32 %r532, %r531, %r521;
add.s32 %r533, %r45, %r505;
add.s32 %r534, %r533, %r532;
add.s32 %r535, %r534, -343485551;
shf.l.wrap.b32 %r536, %r535, %r535, 21;
add.s32 %r537, %r513, 1732584193;
add.s32 %r538, %r529, %r536;
add.s32 %r539, %r538, -271733879;
add.s32 %r540, %r529, -1732584194;
add.s32 %r541, %r521, 271733878;
mul.wide.u32 %rd8, %r1, 16;
add.s64 %rd9, %rd2, %rd8;
st.global.u32 [%rd9], %r537;
st.global.u32 [%rd9+4], %r539;
st.global.u32 [%rd9+8], %r540;
st.global.u32 [%rd9+12], %r541;
BB1_2:
ret;
}
// .globl m00400_loop
.entry m00400_loop(
.param .u64 .ptr .global .align 4 m00400_loop_param_0,
.param .u64 .ptr .global .align 4 m00400_loop_param_1,
.param .u64 .ptr .global .align 4 m00400_loop_param_2,
.param .u64 .ptr .global .align 4 m00400_loop_param_3,
.param .u64 .ptr .global .align 4 m00400_loop_param_4,
.param .u64 .ptr .global .align 1 m00400_loop_param_5,
.param .u64 .ptr .global .align 4 m00400_loop_param_6,
.param .u64 .ptr .global .align 4 m00400_loop_param_7,
.param .u64 .ptr .global .align 4 m00400_loop_param_8,
.param .u64 .ptr .global .align 4 m00400_loop_param_9,
.param .u64 .ptr .global .align 4 m00400_loop_param_10,
.param .u64 .ptr .global .align 4 m00400_loop_param_11,
.param .u64 .ptr .global .align 4 m00400_loop_param_12,
.param .u64 .ptr .global .align 4 m00400_loop_param_13,
.param .u64 .ptr .global .align 4 m00400_loop_param_14,
.param .u64 .ptr .global .align 4 m00400_loop_param_15,
.param .u64 .ptr .global .align 4 m00400_loop_param_16,
.param .u64 .ptr .global .align 4 m00400_loop_param_17,
.param .u64 .ptr .global .align 1 m00400_loop_param_18,
.param .u64 .ptr .global .align 4 m00400_loop_param_19,
.param .u64 .ptr .global .align 4 m00400_loop_param_20,
.param .u64 .ptr .global .align 4 m00400_loop_param_21,
.param .u64 .ptr .global .align 4 m00400_loop_param_22,
.param .u64 .ptr .global .align 4 m00400_loop_param_23,
.param .u32 m00400_loop_param_24,
.param .u32 m00400_loop_param_25,
.param .u32 m00400_loop_param_26,
.param .u32 m00400_loop_param_27,
.param .u32 m00400_loop_param_28,
.param .u32 m00400_loop_param_29,
.param .u32 m00400_loop_param_30,
.param .u32 m00400_loop_param_31,
.param .u32 m00400_loop_param_32,
.param .u32 m00400_loop_param_33,
.param .u32 m00400_loop_param_34
)
{
.reg .pred %p<17>;
.reg .b32 %r<554>;
.reg .b64 %rd<8>;
ld.param.u64 %rd3, [m00400_loop_param_0];
ld.param.u64 %rd4, [m00400_loop_param_4];
ld.param.u32 %r31, [m00400_loop_param_29];
ld.param.u32 %r32, [m00400_loop_param_34];
mov.b32 %r33, %envreg3;
mov.u32 %r34, %ctaid.x;
mov.u32 %r35, %ntid.x;
mad.lo.s32 %r36, %r34, %r35, %r33;
mov.u32 %r37, %tid.x;
add.s32 %r1, %r36, %r37;
setp.ge.u32 %p1, %r1, %r32;
@%p1 bra BB2_5;
cvt.u64.u32 %rd1, %r1;
mul.wide.u32 %rd5, %r1, 16;
add.s64 %rd2, %rd4, %rd5;
ld.global.u32 %r550, [%rd2];
ld.global.u32 %r551, [%rd2+4];
ld.global.u32 %r552, [%rd2+8];
ld.global.u32 %r553, [%rd2+12];
setp.eq.s32 %p2, %r31, 0;
@%p2 bra BB2_4;
mul.lo.s64 %rd6, %rd1, 80;
add.s64 %rd7, %rd3, %rd6;
ld.global.u32 %r39, [%rd7+64];
add.s32 %r40, %r39, 16;
shl.b32 %r41, %r40, 3;
and.b32 %r42, %r41, 24;
mov.u32 %r43, 128;
shl.b32 %r44, %r43, %r42;
setp.lt.u32 %p3, %r39, -16;
setp.lt.u32 %p4, %r40, 20;
and.pred %p5, %p3, %p4;
selp.b32 %r45, %r44, 0, %p5;
ld.global.u32 %r46, [%rd7];
or.b32 %r47, %r45, %r46;
and.b32 %r48, %r40, -4;
setp.eq.s32 %p6, %r48, 20;
selp.b32 %r49, %r44, 0, %p6;
ld.global.u32 %r50, [%rd7+4];
or.b32 %r51, %r49, %r50;
setp.eq.s32 %p7, %r48, 24;
selp.b32 %r52, %r44, 0, %p7;
ld.global.u32 %r53, [%rd7+8];
or.b32 %r54, %r52, %r53;
setp.eq.s32 %p8, %r48, 28;
selp.b32 %r55, %r44, 0, %p8;
ld.global.u32 %r56, [%rd7+12];
or.b32 %r57, %r55, %r56;
setp.eq.s32 %p9, %r48, 32;
selp.b32 %r58, %r44, 0, %p9;
ld.global.u32 %r59, [%rd7+16];
or.b32 %r60, %r58, %r59;
setp.eq.s32 %p10, %r48, 36;
selp.b32 %r61, %r44, 0, %p10;
ld.global.u32 %r62, [%rd7+20];
or.b32 %r63, %r61, %r62;
setp.eq.s32 %p11, %r48, 40;
selp.b32 %r64, %r44, 0, %p11;
ld.global.u32 %r65, [%rd7+24];
or.b32 %r66, %r64, %r65;
setp.eq.s32 %p12, %r48, 44;
selp.b32 %r67, %r44, 0, %p12;
ld.global.u32 %r68, [%rd7+28];
or.b32 %r69, %r67, %r68;
setp.eq.s32 %p13, %r48, 48;
selp.b32 %r70, %r44, 0, %p13;
ld.global.u32 %r71, [%rd7+32];
or.b32 %r72, %r70, %r71;
setp.eq.s32 %p14, %r48, 52;
selp.b32 %r73, %r44, 0, %p14;
ld.global.u32 %r74, [%rd7+36];
or.b32 %r75, %r73, %r74;
setp.eq.s32 %p15, %r48, 56;
selp.b32 %r76, %r44, 0, %p15;
or.b32 %r77, %r76, %r41;
add.s32 %r6, %r47, -448152776;
add.s32 %r7, %r51, 1200080426;
add.s32 %r8, %r54, -1473231341;
add.s32 %r9, %r57, -45705983;
add.s32 %r10, %r60, 1770035416;
add.s32 %r11, %r63, -1958414417;
add.s32 %r12, %r66, -42063;
add.s32 %r13, %r69, -1990404162;
add.s32 %r14, %r72, 1804603682;
add.s32 %r15, %r75, -40341101;
add.s32 %r16, %r77, -1502002290;
mov.u32 %r549, 0;
BB2_3:
add.s32 %r78, %r550, -680876937;
shf.l.wrap.b32 %r79, %r78, %r78, 7;
add.s32 %r80, %r79, -271733879;
and.b32 %r81, %r80, 2004318071;
xor.b32 %r82, %r81, -1732584194;
add.s32 %r83, %r551, %r82;
add.s32 %r84, %r83, -117830708;
shf.l.wrap.b32 %r85, %r84, %r84, 12;
add.s32 %r86, %r85, %r80;
xor.b32 %r87, %r80, -271733879;
and.b32 %r88, %r86, %r87;
xor.b32 %r89, %r88, -271733879;
add.s32 %r90, %r552, %r89;
add.s32 %r91, %r90, -1126478375;
shf.l.wrap.b32 %r92, %r91, %r91, 17;
add.s32 %r93, %r92, %r86;
xor.b32 %r94, %r86, %r80;
and.b32 %r95, %r93, %r94;
xor.b32 %r96, %r95, %r80;
add.s32 %r97, %r553, %r96;
add.s32 %r98, %r97, -1316259209;
shf.l.wrap.b32 %r99, %r98, %r98, 22;
add.s32 %r100, %r99, %r93;
xor.b32 %r101, %r93, %r86;
and.b32 %r102, %r100, %r101;
xor.b32 %r103, %r102, %r86;
add.s32 %r104, %r6, %r79;
add.s32 %r105, %r104, %r103;
shf.l.wrap.b32 %r106, %r105, %r105, 7;
add.s32 %r107, %r106, %r100;
xor.b32 %r108, %r100, %r93;
and.b32 %r109, %r107, %r108;
xor.b32 %r110, %r109, %r93;
add.s32 %r111, %r7, %r86;
add.s32 %r112, %r111, %r110;
shf.l.wrap.b32 %r113, %r112, %r112, 12;
add.s32 %r114, %r113, %r107;
xor.b32 %r115, %r107, %r100;
and.b32 %r116, %r114, %r115;
xor.b32 %r117, %r116, %r100;
add.s32 %r118, %r8, %r93;
add.s32 %r119, %r118, %r117;
shf.l.wrap.b32 %r120, %r119, %r119, 17;
add.s32 %r121, %r120, %r114;
xor.b32 %r122, %r114, %r107;
and.b32 %r123, %r121, %r122;
xor.b32 %r124, %r123, %r107;
add.s32 %r125, %r9, %r100;
add.s32 %r126, %r125, %r124;
shf.l.wrap.b32 %r127, %r126, %r126, 22;
add.s32 %r128, %r127, %r121;
xor.b32 %r129, %r121, %r114;
and.b32 %r130, %r128, %r129;
xor.b32 %r131, %r130, %r114;
add.s32 %r132, %r10, %r107;
add.s32 %r133, %r132, %r131;
shf.l.wrap.b32 %r134, %r133, %r133, 7;
add.s32 %r135, %r134, %r128;
xor.b32 %r136, %r128, %r121;
and.b32 %r137, %r135, %r136;
xor.b32 %r138, %r137, %r121;
add.s32 %r139, %r11, %r114;
add.s32 %r140, %r139, %r138;
shf.l.wrap.b32 %r141, %r140, %r140, 12;
add.s32 %r142, %r141, %r135;
xor.b32 %r143, %r135, %r128;
and.b32 %r144, %r142, %r143;
xor.b32 %r145, %r144, %r128;
add.s32 %r146, %r12, %r121;
add.s32 %r147, %r146, %r145;
shf.l.wrap.b32 %r148, %r147, %r147, 17;
add.s32 %r149, %r148, %r142;
xor.b32 %r150, %r142, %r135;
and.b32 %r151, %r149, %r150;
xor.b32 %r152, %r151, %r135;
add.s32 %r153, %r13, %r128;
add.s32 %r154, %r153, %r152;
shf.l.wrap.b32 %r155, %r154, %r154, 22;
add.s32 %r156, %r155, %r149;
xor.b32 %r157, %r149, %r142;
and.b32 %r158, %r156, %r157;
xor.b32 %r159, %r158, %r142;
add.s32 %r160, %r14, %r135;
add.s32 %r161, %r160, %r159;
shf.l.wrap.b32 %r162, %r161, %r161, 7;
add.s32 %r163, %r162, %r156;
xor.b32 %r164, %r156, %r149;
and.b32 %r165, %r163, %r164;
xor.b32 %r166, %r165, %r149;
add.s32 %r167, %r15, %r142;
add.s32 %r168, %r167, %r166;
shf.l.wrap.b32 %r169, %r168, %r168, 12;
add.s32 %r170, %r169, %r163;
xor.b32 %r171, %r163, %r156;
and.b32 %r172, %r170, %r171;
xor.b32 %r173, %r172, %r156;
add.s32 %r174, %r16, %r149;
add.s32 %r175, %r174, %r173;
shf.l.wrap.b32 %r176, %r175, %r175, 17;
add.s32 %r177, %r176, %r170;
xor.b32 %r178, %r170, %r163;
and.b32 %r179, %r177, %r178;
xor.b32 %r180, %r179, %r163;
add.s32 %r181, %r156, %r180;
add.s32 %r182, %r181, 1236535329;
shf.l.wrap.b32 %r183, %r182, %r182, 22;
add.s32 %r184, %r183, %r177;
xor.b32 %r185, %r184, %r177;
and.b32 %r186, %r185, %r170;
xor.b32 %r187, %r186, %r177;
add.s32 %r188, %r551, %r163;
add.s32 %r189, %r188, %r187;
add.s32 %r190, %r189, -165796510;
shf.l.wrap.b32 %r191, %r190, %r190, 5;
add.s32 %r192, %r191, %r184;
xor.b32 %r193, %r192, %r184;
and.b32 %r194, %r193, %r177;
xor.b32 %r195, %r194, %r184;
add.s32 %r196, %r8, %r170;
add.s32 %r197, %r196, %r195;
add.s32 %r198, %r197, 403729709;
shf.l.wrap.b32 %r199, %r198, %r198, 9;
add.s32 %r200, %r199, %r192;
xor.b32 %r201, %r200, %r192;
and.b32 %r202, %r201, %r184;
xor.b32 %r203, %r202, %r192;
add.s32 %r204, %r13, %r177;
add.s32 %r205, %r204, %r203;
add.s32 %r206, %r205, -1660845421;
shf.l.wrap.b32 %r207, %r206, %r206, 14;
add.s32 %r208, %r207, %r200;
xor.b32 %r209, %r208, %r200;
and.b32 %r210, %r209, %r192;
xor.b32 %r211, %r210, %r200;
add.s32 %r212, %r550, %r184;
add.s32 %r213, %r212, %r211;
add.s32 %r214, %r213, -373897302;
shf.l.wrap.b32 %r215, %r214, %r214, 20;
add.s32 %r216, %r215, %r208;
xor.b32 %r217, %r216, %r208;
and.b32 %r218, %r217, %r200;
xor.b32 %r219, %r218, %r208;
add.s32 %r220, %r7, %r192;
add.s32 %r221, %r220, %r219;
add.s32 %r222, %r221, -1901639117;
shf.l.wrap.b32 %r223, %r222, %r222, 5;
add.s32 %r224, %r223, %r216;
xor.b32 %r225, %r224, %r216;
and.b32 %r226, %r225, %r208;
xor.b32 %r227, %r226, %r216;
add.s32 %r228, %r12, %r200;
add.s32 %r229, %r228, %r227;
add.s32 %r230, %r229, 38058146;
shf.l.wrap.b32 %r231, %r230, %r230, 9;
add.s32 %r232, %r231, %r224;
xor.b32 %r233, %r232, %r224;
and.b32 %r234, %r233, %r216;
xor.b32 %r235, %r234, %r224;
add.s32 %r236, %r208, %r235;
add.s32 %r237, %r236, -660478335;
shf.l.wrap.b32 %r238, %r237, %r237, 14;
add.s32 %r239, %r238, %r232;
xor.b32 %r240, %r239, %r232;
and.b32 %r241, %r240, %r224;
xor.b32 %r242, %r241, %r232;
add.s32 %r243, %r6, %r216;
add.s32 %r244, %r243, %r242;
add.s32 %r245, %r244, 42614928;
shf.l.wrap.b32 %r246, %r245, %r245, 20;
add.s32 %r247, %r246, %r239;
xor.b32 %r248, %r247, %r239;
and.b32 %r249, %r248, %r232;
xor.b32 %r250, %r249, %r239;
add.s32 %r251, %r11, %r224;
add.s32 %r252, %r251, %r250;
add.s32 %r253, %r252, -1768106441;
shf.l.wrap.b32 %r254, %r253, %r253, 5;
add.s32 %r255, %r254, %r247;
xor.b32 %r256, %r255, %r247;
and.b32 %r257, %r256, %r239;
xor.b32 %r258, %r257, %r247;
add.s32 %r259, %r16, %r232;
add.s32 %r260, %r259, %r258;
add.s32 %r261, %r260, 482198600;
shf.l.wrap.b32 %r262, %r261, %r261, 9;
add.s32 %r263, %r262, %r255;
xor.b32 %r264, %r263, %r255;
and.b32 %r265, %r264, %r247;
xor.b32 %r266, %r265, %r255;
add.s32 %r267, %r553, %r239;
add.s32 %r268, %r267, %r266;
add.s32 %r269, %r268, -187363961;
shf.l.wrap.b32 %r270, %r269, %r269, 14;
add.s32 %r271, %r270, %r263;
xor.b32 %r272, %r271, %r263;
and.b32 %r273, %r272, %r255;
xor.b32 %r274, %r273, %r263;
add.s32 %r275, %r10, %r247;
add.s32 %r276, %r275, %r274;
add.s32 %r277, %r276, -606503915;
shf.l.wrap.b32 %r278, %r277, %r277, 20;
add.s32 %r279, %r278, %r271;
xor.b32 %r280, %r279, %r271;
and.b32 %r281, %r280, %r263;
xor.b32 %r282, %r281, %r271;
add.s32 %r283, %r15, %r255;
add.s32 %r284, %r283, %r282;
add.s32 %r285, %r284, -1404340366;
shf.l.wrap.b32 %r286, %r285, %r285, 5;
add.s32 %r287, %r286, %r279;
xor.b32 %r288, %r287, %r279;
and.b32 %r289, %r288, %r271;
xor.b32 %r290, %r289, %r279;
add.s32 %r291, %r552, %r263;
add.s32 %r292, %r291, %r290;
add.s32 %r293, %r292, -51403784;
shf.l.wrap.b32 %r294, %r293, %r293, 9;
add.s32 %r295, %r294, %r287;
xor.b32 %r296, %r295, %r287;
and.b32 %r297, %r296, %r279;
xor.b32 %r298, %r297, %r287;
add.s32 %r299, %r9, %r271;
add.s32 %r300, %r299, %r298;
add.s32 %r301, %r300, 1781034456;
shf.l.wrap.b32 %r302, %r301, %r301, 14;
add.s32 %r303, %r302, %r295;
xor.b32 %r304, %r303, %r295;
and.b32 %r305, %r304, %r287;
xor.b32 %r306, %r305, %r295;
add.s32 %r307, %r14, %r279;
add.s32 %r308, %r307, %r306;
add.s32 %r309, %r308, 563755880;
shf.l.wrap.b32 %r310, %r309, %r309, 20;
add.s32 %r311, %r310, %r303;
xor.b32 %r312, %r304, %r311;
add.s32 %r313, %r7, %r287;
add.s32 %r314, %r313, %r312;
add.s32 %r315, %r314, -1200458984;
shf.l.wrap.b32 %r316, %r315, %r315, 4;
add.s32 %r317, %r316, %r311;
xor.b32 %r318, %r311, %r303;
xor.b32 %r319, %r318, %r317;
add.s32 %r320, %r10, %r295;
add.s32 %r321, %r320, %r319;
add.s32 %r322, %r321, 502357417;
shf.l.wrap.b32 %r323, %r322, %r322, 11;
add.s32 %r324, %r323, %r317;
xor.b32 %r325, %r317, %r311;
xor.b32 %r326, %r325, %r324;
add.s32 %r327, %r13, %r303;
add.s32 %r328, %r327, %r326;
add.s32 %r329, %r328, -465532572;
shf.l.wrap.b32 %r330, %r329, %r329, 16;
add.s32 %r331, %r330, %r324;
xor.b32 %r332, %r324, %r317;
xor.b32 %r333, %r332, %r331;
add.s32 %r334, %r16, %r311;
add.s32 %r335, %r334, %r333;
add.s32 %r336, %r335, 1466692734;
shf.l.wrap.b32 %r337, %r336, %r336, 23;
add.s32 %r338, %r337, %r331;
xor.b32 %r339, %r331, %r324;
xor.b32 %r340, %r339, %r338;
add.s32 %r341, %r551, %r317;
add.s32 %r342, %r341, %r340;
add.s32 %r343, %r342, -1530992060;
shf.l.wrap.b32 %r344, %r343, %r343, 4;
add.s32 %r345, %r344, %r338;
xor.b32 %r346, %r338, %r331;
xor.b32 %r347, %r346, %r345;
add.s32 %r348, %r6, %r324;
add.s32 %r349, %r348, %r347;
add.s32 %r350, %r349, 1721046129;
shf.l.wrap.b32 %r351, %r350, %r350, 11;
add.s32 %r352, %r351, %r345;
xor.b32 %r353, %r345, %r338;
xor.b32 %r354, %r353, %r352;
add.s32 %r355, %r9, %r331;
add.s32 %r356, %r355, %r354;
add.s32 %r357, %r356, -109791649;
shf.l.wrap.b32 %r358, %r357, %r357, 16;
add.s32 %r359, %r358, %r352;
xor.b32 %r360, %r352, %r345;
xor.b32 %r361, %r360, %r359;
add.s32 %r362, %r12, %r338;
add.s32 %r363, %r362, %r361;
add.s32 %r364, %r363, -1094688577;
shf.l.wrap.b32 %r365, %r364, %r364, 23;
add.s32 %r366, %r365, %r359;
xor.b32 %r367, %r359, %r352;
xor.b32 %r368, %r367, %r366;
add.s32 %r369, %r15, %r345;
add.s32 %r370, %r369, %r368;
add.s32 %r371, %r370, 721620275;
shf.l.wrap.b32 %r372, %r371, %r371, 4;
add.s32 %r373, %r372, %r366;
xor.b32 %r374, %r366, %r359;
xor.b32 %r375, %r374, %r373;
add.s32 %r376, %r550, %r352;
add.s32 %r377, %r376, %r375;
add.s32 %r378, %r377, -358537222;
shf.l.wrap.b32 %r379, %r378, %r378, 11;
add.s32 %r380, %r379, %r373;
xor.b32 %r381, %r373, %r366;
xor.b32 %r382, %r381, %r380;
add.s32 %r383, %r553, %r359;
add.s32 %r384, %r383, %r382;
add.s32 %r385, %r384, -722521979;
shf.l.wrap.b32 %r386, %r385, %r385, 16;
add.s32 %r387, %r386, %r380;
xor.b32 %r388, %r380, %r373;
xor.b32 %r389, %r388, %r387;
add.s32 %r390, %r8, %r366;
add.s32 %r391, %r390, %r389;
add.s32 %r392, %r391, 1549260530;
shf.l.wrap.b32 %r393, %r392, %r392, 23;
add.s32 %r394, %r393, %r387;
xor.b32 %r395, %r387, %r380;
xor.b32 %r396, %r395, %r394;
add.s32 %r397, %r11, %r373;
add.s32 %r398, %r397, %r396;
add.s32 %r399, %r398, 1318049930;
shf.l.wrap.b32 %r400, %r399, %r399, 4;
add.s32 %r401, %r400, %r394;
xor.b32 %r402, %r394, %r387;
xor.b32 %r403, %r402, %r401;
add.s32 %r404, %r14, %r380;
add.s32 %r405, %r404, %r403;
add.s32 %r406, %r405, 2068547779;
shf.l.wrap.b32 %r407, %r406, %r406, 11;
add.s32 %r408, %r407, %r401;
xor.b32 %r409, %r401, %r394;
xor.b32 %r410, %r409, %r408;
add.s32 %r411, %r387, %r410;
add.s32 %r412, %r411, 530742520;
shf.l.wrap.b32 %r413, %r412, %r412, 16;
add.s32 %r414, %r413, %r408;
xor.b32 %r415, %r408, %r401;
xor.b32 %r416, %r415, %r414;
add.s32 %r417, %r552, %r394;
add.s32 %r418, %r417, %r416;
add.s32 %r419, %r418, -995338651;
shf.l.wrap.b32 %r420, %r419, %r419, 23;
add.s32 %r421, %r420, %r414;
not.b32 %r422, %r408;
or.b32 %r423, %r421, %r422;
xor.b32 %r424, %r423, %r414;
add.s32 %r425, %r550, %r401;
add.s32 %r426, %r425, %r424;
add.s32 %r427, %r426, -198630844;
shf.l.wrap.b32 %r428, %r427, %r427, 6;
add.s32 %r429, %r428, %r421;
not.b32 %r430, %r414;
or.b32 %r431, %r429, %r430;
xor.b32 %r432, %r431, %r421;
add.s32 %r433, %r9, %r408;
add.s32 %r434, %r433, %r432;
add.s32 %r435, %r434, 1172597398;
shf.l.wrap.b32 %r436, %r435, %r435, 10;
add.s32 %r437, %r436, %r429;
not.b32 %r438, %r421;
or.b32 %r439, %r437, %r438;
xor.b32 %r440, %r439, %r429;
add.s32 %r441, %r16, %r414;
add.s32 %r442, %r441, %r440;
add.s32 %r443, %r442, 85647385;
shf.l.wrap.b32 %r444, %r443, %r443, 15;
add.s32 %r445, %r444, %r437;
not.b32 %r446, %r429;
or.b32 %r447, %r445, %r446;
xor.b32 %r448, %r447, %r437;
add.s32 %r449, %r7, %r421;
add.s32 %r450, %r449, %r448;
add.s32 %r451, %r450, -1257514481;
shf.l.wrap.b32 %r452, %r451, %r451, 21;
add.s32 %r453, %r452, %r445;
not.b32 %r454, %r437;
or.b32 %r455, %r453, %r454;
xor.b32 %r456, %r455, %r445;
add.s32 %r457, %r14, %r429;
add.s32 %r458, %r457, %r456;
add.s32 %r459, %r458, -104118111;
shf.l.wrap.b32 %r460, %r459, %r459, 6;
add.s32 %r461, %r460, %r453;
not.b32 %r462, %r445;
or.b32 %r463, %r461, %r462;
xor.b32 %r464, %r463, %r453;
add.s32 %r465, %r553, %r437;
add.s32 %r466, %r465, %r464;
add.s32 %r467, %r466, -1894986606;
shf.l.wrap.b32 %r468, %r467, %r467, 10;
add.s32 %r469, %r468, %r461;
not.b32 %r470, %r453;
or.b32 %r471, %r469, %r470;
xor.b32 %r472, %r471, %r461;
add.s32 %r473, %r12, %r445;
add.s32 %r474, %r473, %r472;
add.s32 %r475, %r474, -1009460;
shf.l.wrap.b32 %r476, %r475, %r475, 15;
add.s32 %r477, %r476, %r469;
not.b32 %r478, %r461;
or.b32 %r479, %r477, %r478;
xor.b32 %r480, %r479, %r469;
add.s32 %r481, %r551, %r453;
add.s32 %r482, %r481, %r480;
add.s32 %r483, %r482, -2054922799;
shf.l.wrap.b32 %r484, %r483, %r483, 21;
add.s32 %r485, %r484, %r477;
not.b32 %r486, %r469;
or.b32 %r487, %r485, %r486;
xor.b32 %r488, %r487, %r477;
add.s32 %r489, %r10, %r461;
add.s32 %r490, %r489, %r488;
add.s32 %r491, %r490, 103277943;
shf.l.wrap.b32 %r492, %r491, %r491, 6;
add.s32 %r493, %r492, %r485;
not.b32 %r494, %r477;
or.b32 %r495, %r493, %r494;
xor.b32 %r496, %r495, %r485;
add.s32 %r497, %r469, %r496;
add.s32 %r498, %r497, -30611744;
shf.l.wrap.b32 %r499, %r498, %r498, 10;
add.s32 %r500, %r499, %r493;
not.b32 %r501, %r485;
or.b32 %r502, %r500, %r501;
xor.b32 %r503, %r502, %r493;
add.s32 %r504, %r8, %r477;
add.s32 %r505, %r504, %r503;
add.s32 %r506, %r505, -86967039;
shf.l.wrap.b32 %r507, %r506, %r506, 15;
add.s32 %r508, %r507, %r500;
not.b32 %r509, %r493;
or.b32 %r510, %r508, %r509;
xor.b32 %r511, %r510, %r500;
add.s32 %r512, %r15, %r485;
add.s32 %r513, %r512, %r511;
add.s32 %r514, %r513, 1349492750;
shf.l.wrap.b32 %r515, %r514, %r514, 21;
add.s32 %r516, %r515, %r508;
not.b32 %r517, %r500;
or.b32 %r518, %r516, %r517;
xor.b32 %r519, %r518, %r508;
add.s32 %r520, %r6, %r493;
add.s32 %r521, %r520, %r519;
add.s32 %r522, %r521, 302629706;
shf.l.wrap.b32 %r523, %r522, %r522, 6;
add.s32 %r524, %r523, %r516;
not.b32 %r525, %r508;
or.b32 %r526, %r524, %r525;
xor.b32 %r527, %r526, %r516;
add.s32 %r528, %r13, %r500;
add.s32 %r529, %r528, %r527;
add.s32 %r530, %r529, 870193783;
shf.l.wrap.b32 %r531, %r530, %r530, 10;
add.s32 %r532, %r531, %r524;
not.b32 %r533, %r516;
or.b32 %r534, %r532, %r533;
xor.b32 %r535, %r534, %r524;
add.s32 %r536, %r552, %r508;
add.s32 %r537, %r536, %r535;
add.s32 %r538, %r537, 718787259;
shf.l.wrap.b32 %r539, %r538, %r538, 15;
add.s32 %r540, %r539, %r532;
not.b32 %r541, %r524;
or.b32 %r542, %r540, %r541;
xor.b32 %r543, %r542, %r532;
add.s32 %r544, %r11, %r516;
add.s32 %r545, %r544, %r543;
add.s32 %r546, %r545, 1614928866;
shf.l.wrap.b32 %r547, %r546, %r546, 21;
add.s32 %r553, %r532, 271733878;
add.s32 %r552, %r540, -1732584194;
add.s32 %r548, %r540, %r547;
add.s32 %r551, %r548, -271733879;
add.s32 %r550, %r524, 1732584193;
add.s32 %r549, %r549, 1;
setp.lt.u32 %p16, %r549, %r31;
@%p16 bra BB2_3;
BB2_4:
st.global.u32 [%rd2], %r550;
st.global.u32 [%rd2+4], %r551;
st.global.u32 [%rd2+8], %r552;
st.global.u32 [%rd2+12], %r553;
BB2_5:
ret;
}
// .globl m00400_comp
.entry m00400_comp(
.param .u64 .ptr .global .align 4 m00400_comp_param_0,
.param .u64 .ptr .global .align 4 m00400_comp_param_1,
.param .u64 .ptr .global .align 4 m00400_comp_param_2,
.param .u64 .ptr .global .align 4 m00400_comp_param_3,
.param .u64 .ptr .global .align 4 m00400_comp_param_4,
.param .u64 .ptr .global .align 1 m00400_comp_param_5,
.param .u64 .ptr .global .align 4 m00400_comp_param_6,
.param .u64 .ptr .global .align 4 m00400_comp_param_7,
.param .u64 .ptr .global .align 4 m00400_comp_param_8,
.param .u64 .ptr .global .align 4 m00400_comp_param_9,
.param .u64 .ptr .global .align 4 m00400_comp_param_10,
.param .u64 .ptr .global .align 4 m00400_comp_param_11,
.param .u64 .ptr .global .align 4 m00400_comp_param_12,
.param .u64 .ptr .global .align 4 m00400_comp_param_13,
.param .u64 .ptr .global .align 4 m00400_comp_param_14,
.param .u64 .ptr .global .align 4 m00400_comp_param_15,
.param .u64 .ptr .global .align 4 m00400_comp_param_16,
.param .u64 .ptr .global .align 4 m00400_comp_param_17,
.param .u64 .ptr .global .align 1 m00400_comp_param_18,
.param .u64 .ptr .global .align 4 m00400_comp_param_19,
.param .u64 .ptr .global .align 4 m00400_comp_param_20,
.param .u64 .ptr .global .align 4 m00400_comp_param_21,
.param .u64 .ptr .global .align 4 m00400_comp_param_22,
.param .u64 .ptr .global .align 4 m00400_comp_param_23,
.param .u32 m00400_comp_param_24,
.param .u32 m00400_comp_param_25,
.param .u32 m00400_comp_param_26,
.param .u32 m00400_comp_param_27,
.param .u32 m00400_comp_param_28,
.param .u32 m00400_comp_param_29,
.param .u32 m00400_comp_param_30,
.param .u32 m00400_comp_param_31,
.param .u32 m00400_comp_param_32,
.param .u32 m00400_comp_param_33,
.param .u32 m00400_comp_param_34
)
{
.reg .pred %p<25>;
.reg .b32 %r<106>;
.reg .b64 %rd<41>;
ld.param.u64 %rd2, [m00400_comp_param_4];
ld.param.u64 %rd3, [m00400_comp_param_6];
ld.param.u64 %rd4, [m00400_comp_param_7];
ld.param.u64 %rd5, [m00400_comp_param_8];
ld.param.u64 %rd6, [m00400_comp_param_9];
ld.param.u64 %rd7, [m00400_comp_param_10];
ld.param.u64 %rd8, [m00400_comp_param_11];
ld.param.u64 %rd9, [m00400_comp_param_12];
ld.param.u64 %rd10, [m00400_comp_param_13];
ld.param.u64 %rd11, [m00400_comp_param_14];
ld.param.u64 %rd12, [m00400_comp_param_15];
ld.param.u64 %rd13, [m00400_comp_param_16];
ld.param.u64 %rd14, [m00400_comp_param_19];
ld.param.u32 %r27, [m00400_comp_param_24];
ld.param.u32 %r28, [m00400_comp_param_25];
ld.param.u32 %r29, [m00400_comp_param_26];
ld.param.u32 %r30, [m00400_comp_param_27];
ld.param.u32 %r31, [m00400_comp_param_31];
ld.param.u32 %r32, [m00400_comp_param_32];
ld.param.u32 %r33, [m00400_comp_param_34];
mov.b32 %r34, %envreg3;
mov.u32 %r35, %ctaid.x;
mov.u32 %r36, %ntid.x;
mad.lo.s32 %r37, %r35, %r36, %r34;
mov.u32 %r38, %tid.x;
add.s32 %r1, %r37, %r38;
setp.ge.u32 %p1, %r1, %r33;
@%p1 bra BB3_26;
mul.wide.u32 %rd15, %r1, 16;
add.s64 %rd16, %rd2, %rd15;
ld.global.u32 %r2, [%rd16+4];
ld.global.u32 %r3, [%rd16+8];
ld.global.u32 %r4, [%rd16+12];
and.b32 %r5, %r28, 31;
ld.global.u32 %r6, [%rd16];
shr.u32 %r39, %r6, %r5;
and.b32 %r40, %r39, %r27;
mul.wide.u32 %rd17, %r40, 4;
add.s64 %rd18, %rd3, %rd17;
and.b32 %r41, %r6, 31;
mov.u32 %r42, 1;
shl.b32 %r7, %r42, %r41;
ld.global.u32 %r43, [%rd18];
and.b32 %r44, %r43, %r7;
setp.eq.s32 %p2, %r44, 0;
@%p2 bra BB3_26;
shr.u32 %r45, %r2, %r5;
and.b32 %r46, %r45, %r27;
mul.wide.u32 %rd19, %r46, 4;
add.s64 %rd20, %rd4, %rd19;
and.b32 %r47, %r2, 31;
shl.b32 %r8, %r42, %r47;
ld.global.u32 %r49, [%rd20];
and.b32 %r50, %r49, %r8;
setp.eq.s32 %p3, %r50, 0;
@%p3 bra BB3_26;
shr.u32 %r51, %r3, %r5;
and.b32 %r52, %r51, %r27;
mul.wide.u32 %rd21, %r52, 4;
add.s64 %rd22, %rd5, %rd21;
and.b32 %r53, %r3, 31;
shl.b32 %r9, %r42, %r53;
ld.global.u32 %r55, [%rd22];
and.b32 %r56, %r55, %r9;
setp.eq.s32 %p4, %r56, 0;
@%p4 bra BB3_26;
shr.u32 %r57, %r4, %r5;
and.b32 %r58, %r57, %r27;
mul.wide.u32 %rd23, %r58, 4;
add.s64 %rd24, %rd6, %rd23;
and.b32 %r59, %r4, 31;
shl.b32 %r10, %r42, %r59;
ld.global.u32 %r61, [%rd24];
and.b32 %r62, %r61, %r10;
setp.eq.s32 %p5, %r62, 0;
@%p5 bra BB3_26;
and.b32 %r11, %r29, 31;
shr.u32 %r63, %r6, %r11;
and.b32 %r64, %r63, %r27;
mul.wide.u32 %rd25, %r64, 4;
add.s64 %rd26, %rd7, %rd25;
ld.global.u32 %r65, [%rd26];
and.b32 %r66, %r65, %r7;
setp.eq.s32 %p6, %r66, 0;
@%p6 bra BB3_26;
shr.u32 %r67, %r2, %r11;
and.b32 %r68, %r67, %r27;
mul.wide.u32 %rd27, %r68, 4;
add.s64 %rd28, %rd8, %rd27;
ld.global.u32 %r69, [%rd28];
and.b32 %r70, %r69, %r8;
setp.eq.s32 %p7, %r70, 0;
@%p7 bra BB3_26;
shr.u32 %r71, %r3, %r11;
and.b32 %r72, %r71, %r27;
mul.wide.u32 %rd29, %r72, 4;
add.s64 %rd30, %rd9, %rd29;
ld.global.u32 %r73, [%rd30];
and.b32 %r74, %r73, %r9;
setp.eq.s32 %p8, %r74, 0;
@%p8 bra BB3_26;
shr.u32 %r75, %r4, %r11;
and.b32 %r76, %r75, %r27;
mul.wide.u32 %rd31, %r76, 4;
add.s64 %rd32, %rd10, %rd31;
ld.global.u32 %r77, [%rd32];
and.b32 %r78, %r77, %r10;
setp.eq.s32 %p9, %r78, 0;
@%p9 bra BB3_26;
setp.eq.s32 %p10, %r31, 0;
mov.u32 %r97, 0;
mov.u32 %r79, -1;
mov.u32 %r105, %r79;
@%p10 bra BB3_21;
mov.u32 %r96, %r31;
BB3_11:
mov.u32 %r12, %r96;
shr.u32 %r14, %r12, 1;
add.s32 %r15, %r14, %r97;
cvt.u64.u32 %rd33, %r15;
cvt.u64.u32 %rd34, %r32;
add.s64 %rd35, %rd33, %rd34;
shl.b64 %rd36, %rd35, 4;
add.s64 %rd1, %rd12, %rd36;
ld.global.u32 %r16, [%rd1+12];
setp.gt.u32 %p11, %r4, %r16;
mov.u32 %r103, %r42;
@%p11 bra BB3_19;
setp.lt.u32 %p12, %r4, %r16;
mov.u32 %r82, -1;
mov.u32 %r103, %r82;
@%p12 bra BB3_19;
ld.global.u32 %r17, [%rd1+8];
setp.gt.u32 %p13, %r3, %r17;
mov.u32 %r98, %r42;
mov.u32 %r103, %r98;
@%p13 bra BB3_19;
setp.lt.u32 %p14, %r3, %r17;
mov.u32 %r101, %r82;
mov.u32 %r103, %r101;
@%p14 bra BB3_19;
ld.global.u32 %r18, [%rd1+4];
setp.gt.u32 %p15, %r2, %r18;
mov.u32 %r99, %r42;
mov.u32 %r103, %r99;
@%p15 bra BB3_19;
setp.lt.u32 %p16, %r2, %r18;
mov.u32 %r102, %r82;
mov.u32 %r103, %r102;
@%p16 bra BB3_19;
ld.global.u32 %r19, [%rd1];
setp.gt.u32 %p17, %r6, %r19;
mov.u32 %r100, %r42;
mov.u32 %r103, %r100;
@%p17 bra BB3_19;
setp.lt.u32 %p18, %r6, %r19;
selp.b32 %r20, -1, 0, %p18;
mov.u32 %r103, %r20;
BB3_19:
mov.u32 %r21, %r103;
add.s32 %r88, %r14, 1;
setp.gt.s32 %p19, %r21, 0;
selp.b32 %r89, %r88, 0, %p19;
add.s32 %r97, %r89, %r97;
selp.b32 %r90, -1, 0, %p19;
add.s32 %r91, %r90, %r12;
shr.u32 %r23, %r91, 1;
setp.eq.s32 %p20, %r21, 0;
mov.u32 %r105, %r15;
@%p20 bra BB3_21;
setp.ne.s32 %p21, %r23, 0;
mov.u32 %r96, %r23;
mov.u32 %r104, %r79;
mov.u32 %r105, %r104;
@%p21 bra BB3_11;
BB3_21:
setp.eq.s32 %p22, %r105, -1;
@%p22 bra BB3_26;
add.s32 %r25, %r105, %r32;
mul.wide.u32 %rd37, %r25, 4;
add.s64 %rd38, %rd13, %rd37;
atom.global.add.u32 %r93, [%rd38], 1;
setp.ne.s32 %p23, %r93, 0;
@%p23 bra BB3_26;
atom.global.add.u32 %r26, [%rd14], 1;
setp.lt.u32 %p24, %r26, %r31;
@%p24 bra BB3_25;
bra.uni BB3_24;
BB3_25:
mul.wide.u32 %rd39, %r26, 20;
add.s64 %rd40, %rd11, %rd39;
st.global.u32 [%rd40], %r30;
st.global.u32 [%rd40+4], %r105;
st.global.u32 [%rd40+8], %r25;
st.global.u32 [%rd40+12], %r1;
mov.u32 %r95, 0;
st.global.u32 [%rd40+16], %r95;
bra.uni BB3_26;
BB3_24:
atom.global.add.u32 %r94, [%rd14], -1;
BB3_26:
ret;
}
Reference in New Issue View Git Blame Copy Permalink