| // | |
| // Generated by LLVM NVPTX Back-End | |
| // | |
| .version 8.2 | |
| .target sm_89 | |
| .address_size 64 | |
| // .globl triton__0d1d2d3d4d5d6e7de | |
| .extern .func __assertfail | |
| ( | |
| .param .b64 __assertfail_param_0, | |
| .param .b64 __assertfail_param_1, | |
| .param .b32 __assertfail_param_2, | |
| .param .b64 __assertfail_param_3, | |
| .param .b64 __assertfail_param_4 | |
| ) | |
| ; | |
| .global .align 1 .b8 assertFunc_0[25] = {95, 99, 97, 108, 108, 95, 119, 105, 116, 104, 95, 102, 114, 97, 109, 101, 115, 95, 114, 101, 109, 111, 118, 101, 100}; | |
| .global .align 1 .b8 assertFile_0[38] = {60, 102, 114, 111, 122, 101, 110, 32, 105, 109, 112, 111, 114, 116, 108, 105, 98, 46, 95, 98, 111, 111, 116, 115, 116, 114, 97, 112, 95, 101, 120, 116, 101, 114, 110, 97, 108, 62}; | |
| .global .align 1 .b8 assertMessage_0[39] = {105, 110, 100, 101, 120, 32, 111, 117, 116, 32, 111, 102, 32, 98, 111, 117, 110, 100, 115, 58, 32, 48, 32, 60, 61, 32, 116, 109, 112, 49, 49, 32, 60, 32, 53, 48, 50, 53, 55}; | |
| .extern .shared .align 1 .b8 global_smem[]; | |
| .global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0}; | |
| .visible .entry triton__0d1d2d3d4d5d6e7de( | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_0, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_1, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_2, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_3, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_4, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_5, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_6, | |
| .param .u64 triton__0d1d2d3d4d5d6e7de_param_7 | |
| ) | |
| .maxntid 256, 1, 1 | |
| { | |
| .reg .pred %p<145>; | |
| .reg .b16 %rs<83>; | |
| .reg .b32 %r<247>; | |
| .reg .f32 %f<401>; | |
| .reg .b64 %rd<209>; | |
| .loc 1 18 0 | |
| $L__func_begin0: | |
| .loc 1 18 0 | |
| ld.param.u64 %rd40, [triton__0d1d2d3d4d5d6e7de_param_5]; | |
| ld.param.u64 %rd39, [triton__0d1d2d3d4d5d6e7de_param_4]; | |
| ld.param.u64 %rd49, [triton__0d1d2d3d4d5d6e7de_param_0]; | |
| ld.param.u64 %rd50, [triton__0d1d2d3d4d5d6e7de_param_1]; | |
| $L__tmp0: | |
| .loc 1 24 33 | |
| mov.u32 %r1, %tid.x; | |
| ld.param.u64 %rd51, [triton__0d1d2d3d4d5d6e7de_param_2]; | |
| and.b32 %r2, %r1, 255; | |
| ld.param.u64 %rd52, [triton__0d1d2d3d4d5d6e7de_param_3]; | |
| or.b32 %r36, %r2, 256; | |
| or.b32 %r37, %r2, 512; | |
| or.b32 %r38, %r2, 768; | |
| shl.b32 %r3, %r2, 2; | |
| or.b32 %r39, %r3, 1; | |
| or.b32 %r40, %r3, 2; | |
| or.b32 %r41, %r3, 3; | |
| .loc 1 21 28 | |
| mov.u32 %r34, %ctaid.x; | |
| .loc 1 21 34 | |
| cvt.s64.s32 %rd1, %r34; | |
| .loc 1 23 21 | |
| setp.lt.s32 %p1, %r34, 8; | |
| shl.b32 %r42, %r2, 3; | |
| shl.b32 %r43, %r2, 5; | |
| mov.u32 %r44, global_smem; | |
| add.s32 %r4, %r44, %r43; | |
| shl.b32 %r45, %r39, 1; | |
| shl.b32 %r46, %r39, 3; | |
| add.s32 %r5, %r44, %r46; | |
| shl.b32 %r47, %r40, 1; | |
| shl.b32 %r48, %r40, 3; | |
| add.s32 %r6, %r44, %r48; | |
| shl.b32 %r49, %r41, 1; | |
| shl.b32 %r50, %r41, 3; | |
| add.s32 %r7, %r44, %r50; | |
| shl.b32 %r51, %r2, 1; | |
| add.s32 %r12, %r44, %r42; | |
| shl.b32 %r52, %r36, 1; | |
| shl.b32 %r53, %r36, 3; | |
| add.s32 %r9, %r44, %r53; | |
| shl.b32 %r54, %r37, 1; | |
| shl.b32 %r55, %r37, 3; | |
| add.s32 %r10, %r44, %r55; | |
| shl.b32 %r56, %r38, 1; | |
| shl.b32 %r57, %r38, 3; | |
| add.s32 %r11, %r44, %r57; | |
| add.s32 %r13, %r44, %r45; | |
| add.s32 %r14, %r44, %r47; | |
| add.s32 %r15, %r44, %r49; | |
| add.s32 %r16, %r44, %r51; | |
| add.s32 %r17, %r44, %r52; | |
| add.s32 %r18, %r44, %r54; | |
| add.s32 %r19, %r44, %r56; | |
| add.s32 %r20, %r44, %r3; | |
| add.s32 %r21, %r44, %r2; | |
| shl.b32 %r58, %r2, 6; | |
| add.s32 %r22, %r44, %r58; | |
| shl.b32 %r59, %r39, 4; | |
| add.s32 %r23, %r44, %r59; | |
| shl.b32 %r60, %r40, 4; | |
| add.s32 %r24, %r44, %r60; | |
| shl.b32 %r61, %r41, 4; | |
| add.s32 %r25, %r44, %r61; | |
| shl.b32 %r62, %r2, 4; | |
| add.s32 %r26, %r44, %r62; | |
| shl.b32 %r63, %r36, 4; | |
| add.s32 %r27, %r44, %r63; | |
| shl.b32 %r64, %r37, 4; | |
| add.s32 %r28, %r44, %r64; | |
| shl.b32 %r65, %r38, 4; | |
| add.s32 %r29, %r44, %r65; | |
| .loc 1 28 36 | |
| mul.wide.s32 %rd53, %r34, 61440; | |
| mul.wide.u32 %rd54, %r2, 32; | |
| add.s64 %rd55, %rd53, %rd54; | |
| add.s64 %rd56, %rd55, %rd49; | |
| add.s64 %rd200, %rd56, 8208; | |
| mul.wide.s32 %rd57, %r34, 771947520; | |
| add.s64 %rd58, %rd50, %rd57; | |
| mul.wide.u32 %rd59, %r2, 402056; | |
| add.s64 %rd60, %rd58, %rd59; | |
| add.s64 %rd199, %rd60, 103227878; | |
| mul.wide.u32 %rd4, %r2, 16; | |
| mul.wide.s32 %rd61, %r34, 30720; | |
| add.s64 %rd198, %rd52, %rd61; | |
| add.s64 %rd197, %rd51, %rd61; | |
| mov.u64 %rd201, 0; | |
| mov.f32 %f385, 0f00000000; | |
| mov.b32 %r246, -2048; | |
| mov.u16 %rs44, 0; | |
| mov.f32 %f386, %f385; | |
| mov.f32 %f387, %f385; | |
| mov.f32 %f388, %f385; | |
| mov.f32 %f389, %f385; | |
| mov.f32 %f390, %f385; | |
| mov.f32 %f391, %f385; | |
| mov.f32 %f392, %f385; | |
| mov.u64 %rd202, %rd201; | |
| mov.u64 %rd203, %rd201; | |
| mov.u64 %rd204, %rd201; | |
| mov.u64 %rd205, %rd201; | |
| mov.u64 %rd206, %rd201; | |
| mov.u64 %rd207, %rd201; | |
| mov.u64 %rd208, %rd201; | |
| bra.uni $L__BB0_1; | |
| $L__BB0_19: | |
| .loc 1 0 0 | |
| add.s64 %rd206, %rd206, %rd81; | |
| add.s64 %rd205, %rd205, %rd80; | |
| add.s64 %rd208, %rd208, %rd79; | |
| add.s64 %rd207, %rd207, %rd78; | |
| add.s64 %rd202, %rd202, %rd85; | |
| add.s64 %rd201, %rd201, %rd84; | |
| add.s64 %rd204, %rd204, %rd83; | |
| add.s64 %rd203, %rd203, %rd82; | |
| .loc 1 36 23 | |
| bfe.s32 %r172, %r115, 0, 8; | |
| cvt.u16.u32 %rs67, %r172; | |
| and.b16 %rs68, %rs67, 255; | |
| setp.eq.s16 %p124, %rs68, 0; | |
| bfe.s32 %r173, %r115, 8, 8; | |
| cvt.u16.u32 %rs69, %r173; | |
| and.b16 %rs70, %rs69, 255; | |
| setp.eq.s16 %p125, %rs70, 0; | |
| bfe.s32 %r174, %r115, 16, 8; | |
| cvt.u16.u32 %rs71, %r174; | |
| and.b16 %rs72, %rs71, 255; | |
| setp.eq.s16 %p126, %rs72, 0; | |
| bfe.s32 %r175, %r115, 24, 8; | |
| cvt.u16.u32 %rs73, %r175; | |
| and.b16 %rs74, %rs73, 255; | |
| setp.eq.s16 %p127, %rs74, 0; | |
| bfe.s32 %r176, %r108, 0, 8; | |
| cvt.u16.u32 %rs75, %r176; | |
| and.b16 %rs76, %rs75, 255; | |
| setp.eq.s16 %p128, %rs76, 0; | |
| bfe.s32 %r177, %r108, 8, 8; | |
| cvt.u16.u32 %rs77, %r177; | |
| and.b16 %rs78, %rs77, 255; | |
| setp.eq.s16 %p129, %rs78, 0; | |
| bfe.s32 %r178, %r108, 16, 8; | |
| cvt.u16.u32 %rs79, %r178; | |
| and.b16 %rs80, %rs79, 255; | |
| setp.eq.s16 %p130, %rs80, 0; | |
| bfe.s32 %r179, %r108, 24, 8; | |
| cvt.u16.u32 %rs81, %r179; | |
| and.b16 %rs82, %rs81, 255; | |
| setp.eq.s16 %p131, %rs82, 0; | |
| .loc 1 50 23 | |
| setp.eq.f32 %p132, %f68, 0f00000000; | |
| selp.f32 %f320, 0fFF800000, %f400, %p132; | |
| bar.sync 0; | |
| st.shared.f32 [%r4], %f37; | |
| st.shared.f32 [%r5], %f42; | |
| st.shared.f32 [%r6], %f47; | |
| st.shared.f32 [%r7], %f52; | |
| bar.sync 0; | |
| ld.shared.f32 %f321, [%r12]; | |
| ld.shared.f32 %f322, [%r9]; | |
| ld.shared.f32 %f323, [%r10]; | |
| ld.shared.f32 %f324, [%r11]; | |
| bar.sync 0; | |
| st.shared.f32 [%r4], %f57; | |
| st.shared.f32 [%r5], %f62; | |
| st.shared.f32 [%r6], %f67; | |
| st.shared.f32 [%r7], %f320; | |
| bar.sync 0; | |
| ld.shared.f32 %f325, [%r12]; | |
| ld.shared.f32 %f326, [%r9]; | |
| ld.shared.f32 %f327, [%r10]; | |
| ld.shared.f32 %f328, [%r11]; | |
| .loc 1 54 17 | |
| sub.f32 %f329, %f324, %f28; | |
| sub.f32 %f330, %f323, %f27; | |
| sub.f32 %f331, %f322, %f26; | |
| sub.f32 %f332, %f321, %f25; | |
| sub.f32 %f333, %f328, %f32; | |
| sub.f32 %f334, %f327, %f31; | |
| sub.f32 %f335, %f326, %f30; | |
| sub.f32 %f336, %f325, %f29; | |
| add.f32 %f337, %f336, 0f00000000; | |
| add.f32 %f338, %f335, 0f00000000; | |
| add.f32 %f339, %f334, 0f00000000; | |
| add.f32 %f340, %f333, 0f00000000; | |
| add.f32 %f341, %f332, 0f00000000; | |
| add.f32 %f342, %f331, 0f00000000; | |
| add.f32 %f343, %f330, 0f00000000; | |
| add.f32 %f344, %f329, 0f00000000; | |
| .loc 1 56 38 | |
| selp.f32 %f345, 0f00000000, %f344, %p131; | |
| selp.f32 %f346, 0f00000000, %f343, %p130; | |
| selp.f32 %f347, 0f00000000, %f342, %p129; | |
| selp.f32 %f348, 0f00000000, %f341, %p128; | |
| selp.f32 %f349, 0f00000000, %f340, %p127; | |
| selp.f32 %f350, 0f00000000, %f339, %p126; | |
| selp.f32 %f351, 0f00000000, %f338, %p125; | |
| selp.f32 %f352, 0f00000000, %f337, %p124; | |
| .loc 1 59 48 | |
| selp.f32 %f353, %f352, 0f80000000, %p1; | |
| selp.f32 %f354, %f351, 0f80000000, %p1; | |
| selp.f32 %f355, %f350, 0f80000000, %p97; | |
| selp.f32 %f356, %f349, 0f80000000, %p97; | |
| selp.f32 %f357, %f348, 0f80000000, %p1; | |
| selp.f32 %f358, %f347, 0f80000000, %p1; | |
| selp.f32 %f359, %f346, 0f80000000, %p1; | |
| selp.f32 %f360, %f345, 0f80000000, %p1; | |
| add.f32 %f388, %f388, %f360; | |
| add.f32 %f387, %f387, %f359; | |
| add.f32 %f386, %f386, %f358; | |
| add.f32 %f385, %f385, %f357; | |
| add.f32 %f392, %f392, %f356; | |
| add.f32 %f391, %f391, %f355; | |
| add.f32 %f390, %f390, %f354; | |
| add.f32 %f389, %f389, %f353; | |
| .loc 1 28 36 | |
| add.s64 %rd200, %rd200, 16384; | |
| add.s32 %r246, %r246, 2048; | |
| add.s64 %rd199, %rd199, 205852672; | |
| add.s64 %rd198, %rd198, 8192; | |
| add.s64 %rd197, %rd197, 8192; | |
| setp.lt.u32 %p133, %r246, 5632; | |
| @%p133 bra $L__BB0_1; | |
| bra.uni $L__BB0_20; | |
| $L__BB0_1: | |
| .loc 1 0 36 | |
| cvt.u32.u64 %r98, %rd1; | |
| .loc 1 23 21 | |
| setp.lt.s32 %p85, %r98, 8; | |
| .loc 1 29 27 | |
| add.s32 %r99, %r246, 3584; | |
| add.s32 %r100, %r3, %r246; | |
| add.s32 %r101, %r100, 3072; | |
| .loc 1 30 25 | |
| setp.lt.u32 %p42, %r99, 7680; | |
| setp.lt.u32 %p43, %r101, 7680; | |
| .loc 1 29 27 | |
| add.s64 %rd64, %rd200, -8208; | |
| .loc 1 32 34 | |
| add.s64 %rd67, %rd200, -8192; | |
| add.s64 %rd70, %rd200, -16; | |
| .loc 1 32 59 | |
| and.pred %p97, %p85, %p42; | |
| and.pred %p16, %p85, %p43; | |
| .loc 1 32 51 | |
| mov.u64 %rd62, 0x0; | |
| mov.u64 %rd63, 0x0; | |
| @%p85 ld.global.L1::evict_first.v2.b64 { %rd62, %rd63 }, [ %rd64 + 0 ]; | |
| @!%p85 mov.u64 %rd62, 0x0; | |
| @!%p85 mov.u64 %rd63, 0x0; | |
| mov.u64 %rd65, 0x0; | |
| mov.u64 %rd66, 0x0; | |
| @%p85 ld.global.L1::evict_first.v2.b64 { %rd65, %rd66 }, [ %rd67 + 0 ]; | |
| @!%p85 mov.u64 %rd65, 0x0; | |
| @!%p85 mov.u64 %rd66, 0x0; | |
| mov.u64 %rd68, 0x0; | |
| mov.u64 %rd69, 0x0; | |
| @%p16 ld.global.L1::evict_first.v2.b64 { %rd68, %rd69 }, [ %rd70 + 0 ]; | |
| @!%p16 mov.u64 %rd68, 0x0; | |
| @!%p16 mov.u64 %rd69, 0x0; | |
| mov.u64 %rd71, 0x0; | |
| mov.u64 %rd72, 0x0; | |
| @%p16 ld.global.L1::evict_first.v2.b64 { %rd71, %rd72 }, [ %rd200 + 0 ]; | |
| @!%p16 mov.u64 %rd71, 0x0; | |
| @!%p16 mov.u64 %rd72, 0x0; | |
| .loc 1 33 35 | |
| add.s64 %rd74, %rd197, %rd4; | |
| .loc 1 33 52 | |
| add.s64 %rd75, %rd74, 4096; | |
| mov.b32 %r70, 0; | |
| mov.u32 %r66, 0x0; | |
| mov.u32 %r67, 0x0; | |
| mov.u32 %r68, 0x0; | |
| mov.u32 %r69, 0x0; | |
| @%p85 ld.global.L1::evict_first.v4.b32 { %r66, %r67, %r68, %r69 }, [ %rd74 + 0 ]; | |
| @!%p85 mov.u32 %r66, %r70; | |
| @!%p85 mov.u32 %r67, %r70; | |
| @!%p85 mov.u32 %r68, %r70; | |
| @!%p85 mov.u32 %r69, %r70; | |
| mov.u32 %r74, 0x0; | |
| mov.u32 %r75, 0x0; | |
| mov.u32 %r76, 0x0; | |
| mov.u32 %r77, 0x0; | |
| @%p16 ld.global.L1::evict_first.v4.b32 { %r74, %r75, %r76, %r77 }, [ %rd75 + 0 ]; | |
| @!%p16 mov.u32 %r74, %r70; | |
| @!%p16 mov.u32 %r75, %r70; | |
| @!%p16 mov.u32 %r76, %r70; | |
| @!%p16 mov.u32 %r77, %r70; | |
| bar.sync 0; | |
| st.shared.u32 [%r4], %r66; | |
| st.shared.u32 [%r5], %r67; | |
| st.shared.u32 [%r6], %r68; | |
| st.shared.u32 [%r7], %r69; | |
| bar.sync 0; | |
| ld.shared.f32 %f9, [%r12]; | |
| ld.shared.f32 %f10, [%r9]; | |
| ld.shared.f32 %f11, [%r10]; | |
| ld.shared.f32 %f12, [%r11]; | |
| bar.sync 0; | |
| st.shared.u32 [%r4], %r74; | |
| st.shared.u32 [%r5], %r75; | |
| st.shared.u32 [%r6], %r76; | |
| st.shared.u32 [%r7], %r77; | |
| bar.sync 0; | |
| ld.shared.f32 %f13, [%r12]; | |
| ld.shared.f32 %f14, [%r9]; | |
| ld.shared.f32 %f15, [%r10]; | |
| ld.shared.f32 %f16, [%r11]; | |
| .loc 1 34 35 | |
| add.s64 %rd76, %rd198, %rd4; | |
| .loc 1 34 52 | |
| add.s64 %rd77, %rd76, 4096; | |
| mov.u32 %r82, 0x0; | |
| mov.u32 %r83, 0x0; | |
| mov.u32 %r84, 0x0; | |
| mov.u32 %r85, 0x0; | |
| @%p85 ld.global.L1::evict_first.v4.b32 { %r82, %r83, %r84, %r85 }, [ %rd76 + 0 ]; | |
| @!%p85 mov.u32 %r82, %r70; | |
| @!%p85 mov.u32 %r83, %r70; | |
| @!%p85 mov.u32 %r84, %r70; | |
| @!%p85 mov.u32 %r85, %r70; | |
| mov.b32 %f17, %r82; | |
| mov.u32 %r90, 0x0; | |
| mov.u32 %r91, 0x0; | |
| mov.u32 %r92, 0x0; | |
| mov.u32 %r93, 0x0; | |
| @%p16 ld.global.L1::evict_first.v4.b32 { %r90, %r91, %r92, %r93 }, [ %rd77 + 0 ]; | |
| @!%p16 mov.u32 %r90, %r70; | |
| @!%p16 mov.u32 %r91, %r70; | |
| @!%p16 mov.u32 %r92, %r70; | |
| @!%p16 mov.u32 %r93, %r70; | |
| .loc 1 36 23 | |
| bar.sync 0; | |
| setp.ne.s64 %p44, %rd72, -1; | |
| setp.ne.s64 %p45, %rd71, -1; | |
| setp.ne.s64 %p46, %rd69, -1; | |
| setp.ne.s64 %p47, %rd68, -1; | |
| setp.ne.s64 %p48, %rd66, -1; | |
| setp.ne.s64 %p49, %rd65, -1; | |
| setp.ne.s64 %p50, %rd63, -1; | |
| setp.ne.s64 %p51, %rd62, -1; | |
| selp.u16 %rs1, 1, 0, %p51; | |
| st.shared.u8 [%r12], %rs1; | |
| selp.u16 %rs2, 1, 0, %p50; | |
| st.shared.u8 [%r13], %rs2; | |
| selp.u16 %rs3, 1, 0, %p49; | |
| st.shared.u8 [%r14], %rs3; | |
| selp.u16 %rs4, 1, 0, %p48; | |
| st.shared.u8 [%r15], %rs4; | |
| bar.sync 0; | |
| ld.shared.u8 %r102, [%r19]; | |
| ld.shared.u8 %r103, [%r18]; | |
| ld.shared.u8 %r104, [%r17]; | |
| ld.shared.u8 %r105, [%r16]; | |
| bar.sync 0; | |
| selp.u16 %rs5, 1, 0, %p47; | |
| st.shared.u8 [%r12], %rs5; | |
| selp.u16 %rs6, 1, 0, %p46; | |
| st.shared.u8 [%r13], %rs6; | |
| selp.u16 %rs7, 1, 0, %p45; | |
| st.shared.u8 [%r14], %rs7; | |
| selp.u16 %rs8, 1, 0, %p44; | |
| st.shared.u8 [%r15], %rs8; | |
| bar.sync 0; | |
| bfi.b32 %r106, %r104, %r105, 8, 8; | |
| bfi.b32 %r107, %r103, %r106, 16, 8; | |
| bfi.b32 %r108, %r102, %r107, 24, 8; | |
| ld.shared.u8 %r109, [%r16]; | |
| ld.shared.u8 %r110, [%r17]; | |
| bfi.b32 %r111, %r110, %r109, 8, 8; | |
| ld.shared.u8 %r112, [%r18]; | |
| bfi.b32 %r113, %r112, %r111, 16, 8; | |
| ld.shared.u8 %r114, [%r19]; | |
| bfi.b32 %r115, %r114, %r113, 24, 8; | |
| .loc 1 46 41 | |
| bar.sync 0; | |
| .loc 1 42 36 | |
| selp.b64 %rd86, %rd62, 0, %p51; | |
| selp.b64 %rd87, %rd63, 0, %p50; | |
| selp.b64 %rd88, %rd65, 0, %p49; | |
| selp.b64 %rd89, %rd66, 0, %p48; | |
| .loc 1 43 22 | |
| add.s64 %rd90, %rd89, 50257; | |
| add.s64 %rd91, %rd88, 50257; | |
| add.s64 %rd92, %rd87, 50257; | |
| add.s64 %rd93, %rd86, 50257; | |
| .loc 1 44 23 | |
| setp.lt.s64 %p60, %rd89, 0; | |
| setp.lt.s64 %p61, %rd88, 0; | |
| setp.lt.s64 %p62, %rd87, 0; | |
| setp.lt.s64 %p63, %rd86, 0; | |
| .loc 1 45 38 | |
| selp.b64 %rd27, %rd93, %rd86, %p63; | |
| selp.b64 %rd28, %rd92, %rd87, %p62; | |
| selp.b64 %rd29, %rd91, %rd88, %p61; | |
| selp.b64 %rd30, %rd90, %rd89, %p60; | |
| .loc 1 46 41 | |
| setp.lt.u64 %p64, %rd30, 50257; | |
| setp.lt.u64 %p65, %rd29, 50257; | |
| setp.lt.u64 %p66, %rd28, 50257; | |
| setp.lt.u64 %p67, %rd27, 50257; | |
| selp.u32 %r116, 1, 0, %p67; | |
| selp.u32 %r117, 1, 0, %p66; | |
| bfi.b32 %r118, %r117, %r116, 8, 8; | |
| selp.u32 %r119, 1, 0, %p65; | |
| bfi.b32 %r120, %r119, %r118, 16, 8; | |
| selp.u32 %r121, 1, 0, %p64; | |
| bfi.b32 %r122, %r121, %r120, 24, 8; | |
| st.shared.u32 [%r20], %r122; | |
| bar.sync 0; | |
| ld.shared.u8 %rs9, [%r21]; | |
| ld.shared.u8 %rs10, [%r21+256]; | |
| ld.shared.u8 %rs11, [%r21+512]; | |
| ld.shared.u8 %rs12, [%r21+768]; | |
| bar.sync 0; | |
| .loc 1 42 36 | |
| selp.b64 %rd94, %rd68, 0, %p47; | |
| selp.b64 %rd95, %rd69, 0, %p46; | |
| selp.b64 %rd96, %rd71, 0, %p45; | |
| selp.b64 %rd97, %rd72, 0, %p44; | |
| .loc 1 43 22 | |
| add.s64 %rd98, %rd97, 50257; | |
| add.s64 %rd99, %rd96, 50257; | |
| add.s64 %rd100, %rd95, 50257; | |
| add.s64 %rd101, %rd94, 50257; | |
| .loc 1 44 23 | |
| setp.lt.s64 %p68, %rd97, 0; | |
| setp.lt.s64 %p69, %rd96, 0; | |
| setp.lt.s64 %p70, %rd95, 0; | |
| setp.lt.s64 %p71, %rd94, 0; | |
| .loc 1 45 38 | |
| selp.b64 %rd31, %rd101, %rd94, %p71; | |
| selp.b64 %rd32, %rd100, %rd95, %p70; | |
| selp.b64 %rd33, %rd99, %rd96, %p69; | |
| selp.b64 %rd34, %rd98, %rd97, %p68; | |
| .loc 1 46 41 | |
| setp.lt.u64 %p72, %rd34, 50257; | |
| setp.lt.u64 %p73, %rd33, 50257; | |
| setp.lt.u64 %p74, %rd32, 50257; | |
| setp.lt.u64 %p75, %rd31, 50257; | |
| selp.u32 %r123, 1, 0, %p75; | |
| selp.u32 %r124, 1, 0, %p74; | |
| bfi.b32 %r125, %r124, %r123, 8, 8; | |
| selp.u32 %r126, 1, 0, %p73; | |
| bfi.b32 %r127, %r126, %r125, 16, 8; | |
| selp.u32 %r128, 1, 0, %p72; | |
| bfi.b32 %r129, %r128, %r127, 24, 8; | |
| st.shared.u32 [%r20], %r129; | |
| bar.sync 0; | |
| ld.shared.u8 %rs13, [%r21]; | |
| ld.shared.u8 %rs14, [%r21+256]; | |
| ld.shared.u8 %rs15, [%r21+512]; | |
| ld.shared.u8 %rs16, [%r21+768]; | |
| setp.eq.s16 %p76, %rs11, 0; | |
| selp.u16 %rs17, 1, 0, %p76; | |
| shl.b16 %rs18, %rs17, 2; | |
| setp.eq.s16 %p77, %rs12, 0; | |
| selp.u16 %rs19, -1, 0, %p77; | |
| shl.b16 %rs20, %rs19, 3; | |
| or.b16 %rs21, %rs20, %rs18; | |
| setp.eq.s16 %p78, %rs10, 0; | |
| selp.u16 %rs22, 1, 0, %p78; | |
| setp.eq.s16 %p79, %rs9, 0; | |
| selp.u16 %rs23, -1, 0, %p79; | |
| shl.b16 %rs24, %rs23, 1; | |
| or.b16 %rs25, %rs22, %rs24; | |
| and.b16 %rs26, %rs25, 3; | |
| or.b16 %rs27, %rs26, %rs21; | |
| and.b16 %rs28, %rs27, 15; | |
| setp.eq.s16 %p80, %rs15, 0; | |
| selp.u16 %rs29, 1, 0, %p80; | |
| shl.b16 %rs30, %rs29, 2; | |
| setp.eq.s16 %p81, %rs16, 0; | |
| selp.u16 %rs31, -1, 0, %p81; | |
| shl.b16 %rs32, %rs31, 3; | |
| or.b16 %rs33, %rs32, %rs30; | |
| setp.eq.s16 %p82, %rs13, 0; | |
| selp.u16 %rs34, 1, 0, %p82; | |
| setp.eq.s16 %p83, %rs14, 0; | |
| selp.u16 %rs35, -1, 0, %p83; | |
| shl.b16 %rs36, %rs35, 1; | |
| or.b16 %rs37, %rs34, %rs36; | |
| and.b16 %rs38, %rs37, 3; | |
| or.b16 %rs39, %rs38, %rs33; | |
| shl.b16 %rs40, %rs39, 4; | |
| or.b16 %rs41, %rs28, %rs40; | |
| .loc 1 46 57 | |
| and.b16 %rs42, %rs41, 255; | |
| setp.eq.s16 %p84, %rs42, 0; | |
| @%p84 bra $L__BB0_3; | |
| mov.u64 %rd102, assertMessage_0; | |
| cvta.global.u64 %rd103, %rd102; | |
| mov.u64 %rd104, assertFile_0; | |
| cvta.global.u64 %rd105, %rd104; | |
| mov.u64 %rd106, assertFunc_0; | |
| cvta.global.u64 %rd107, %rd106; | |
| mov.b32 %r130, 883; | |
| mov.u64 %rd108, 1; | |
| { // callseq 0, 0 | |
| .reg .b32 temp_param_reg; | |
| .param .b64 param0; | |
| st.param.b64 [param0+0], %rd103; | |
| .param .b64 param1; | |
| st.param.b64 [param1+0], %rd105; | |
| .param .b32 param2; | |
| st.param.b32 [param2+0], %r130; | |
| .param .b64 param3; | |
| st.param.b64 [param3+0], %rd107; | |
| .param .b64 param4; | |
| st.param.b64 [param4+0], %rd108; | |
| call.uni | |
| __assertfail, | |
| ( | |
| param0, | |
| param1, | |
| param2, | |
| param3, | |
| param4 | |
| ); | |
| } // callseq 0 | |
| $L__BB0_3: | |
| .loc 1 47 73 | |
| bar.sync 0; | |
| shl.b64 %rd117, %rd27, 1; | |
| add.s64 %rd118, %rd199, %rd117; | |
| add.s64 %rd119, %rd118, -103227878; | |
| st.shared.u64 [%r22], %rd119; | |
| shl.b64 %rd120, %rd28, 1; | |
| add.s64 %rd121, %rd199, %rd120; | |
| add.s64 %rd122, %rd121, -103127364; | |
| st.shared.u64 [%r23], %rd122; | |
| shl.b64 %rd123, %rd29, 1; | |
| add.s64 %rd124, %rd199, %rd123; | |
| add.s64 %rd125, %rd124, -103026850; | |
| st.shared.u64 [%r24], %rd125; | |
| shl.b64 %rd126, %rd30, 1; | |
| add.s64 %rd127, %rd199, %rd126; | |
| add.s64 %rd128, %rd127, -102926336; | |
| st.shared.u64 [%r25], %rd128; | |
| bar.sync 0; | |
| ld.shared.u64 %rd109, [%r26]; | |
| ld.shared.u64 %rd110, [%r27]; | |
| ld.shared.u64 %rd111, [%r28]; | |
| ld.shared.u64 %rd112, [%r29]; | |
| bar.sync 0; | |
| shl.b64 %rd129, %rd31, 1; | |
| add.s64 %rd130, %rd199, %rd129; | |
| add.s64 %rd131, %rd130, -301542; | |
| st.shared.u64 [%r22], %rd131; | |
| shl.b64 %rd132, %rd32, 1; | |
| add.s64 %rd133, %rd199, %rd132; | |
| add.s64 %rd134, %rd133, -201028; | |
| st.shared.u64 [%r23], %rd134; | |
| shl.b64 %rd135, %rd33, 1; | |
| add.s64 %rd136, %rd199, %rd135; | |
| add.s64 %rd137, %rd136, -100514; | |
| st.shared.u64 [%r24], %rd137; | |
| shl.b64 %rd138, %rd34, 1; | |
| add.s64 %rd139, %rd199, %rd138; | |
| st.shared.u64 [%r25], %rd139; | |
| bar.sync 0; | |
| ld.shared.u64 %rd113, [%r26]; | |
| ld.shared.u64 %rd114, [%r27]; | |
| ld.shared.u64 %rd115, [%r28]; | |
| ld.shared.u64 %rd116, [%r29]; | |
| mov.u16 %rs43, 0x0; | |
| @%p85 ld.global.L1::evict_last.b16 { %rs43 }, [ %rd109 + 0 ]; | |
| @!%p85 mov.u16 %rs43, %rs44; | |
| mov.u16 %rs45, 0x0; | |
| @%p85 ld.global.L1::evict_last.b16 { %rs45 }, [ %rd110 + 0 ]; | |
| @!%p85 mov.u16 %rs45, %rs44; | |
| mov.u16 %rs47, 0x0; | |
| @%p85 ld.global.L1::evict_last.b16 { %rs47 }, [ %rd111 + 0 ]; | |
| @!%p85 mov.u16 %rs47, %rs44; | |
| mov.u16 %rs49, 0x0; | |
| @%p85 ld.global.L1::evict_last.b16 { %rs49 }, [ %rd112 + 0 ]; | |
| @!%p85 mov.u16 %rs49, %rs44; | |
| mov.u16 %rs51, 0x0; | |
| @%p85 ld.global.L1::evict_last.b16 { %rs51 }, [ %rd113 + 0 ]; | |
| @!%p85 mov.u16 %rs51, %rs44; | |
| mov.u16 %rs53, 0x0; | |
| @%p85 ld.global.L1::evict_last.b16 { %rs53 }, [ %rd114 + 0 ]; | |
| @!%p85 mov.u16 %rs53, %rs44; | |
| mov.u16 %rs55, 0x0; | |
| @%p97 ld.global.L1::evict_last.b16 { %rs55 }, [ %rd115 + 0 ]; | |
| @!%p97 mov.u16 %rs55, %rs44; | |
| mov.u16 %rs57, 0x0; | |
| @%p97 ld.global.L1::evict_last.b16 { %rs57 }, [ %rd116 + 0 ]; | |
| @!%p97 mov.u16 %rs57, %rs44; | |
| .loc 1 50 23 | |
| setp.lt.f32 %p101, %f17, 0f00800000; | |
| mul.f32 %f96, %f17, 0f4B000000; | |
| selp.f32 %f33, %f96, %f17, %p101; | |
| selp.f32 %f97, 0fC1B80000, 0f00000000, %p101; | |
| mov.b32 %r140, %f33; | |
| add.s32 %r141, %r140, -1059760811; | |
| and.b32 %r142, %r141, -8388608; | |
| sub.s32 %r143, %r140, %r142; | |
| mov.b32 %f98, %r143; | |
| cvt.rn.f32.s32 %f99, %r142; | |
| mov.f32 %f100, 0f34000000; | |
| fma.rn.ftz.f32 %f101, %f99, %f100, %f97; | |
| add.f32 %f102, %f98, 0fBF800000; | |
| mov.f32 %f103, 0f3E1039F6; | |
| mov.f32 %f104, 0fBE055027; | |
| fma.rn.ftz.f32 %f105, %f104, %f102, %f103; | |
| mov.f32 %f106, 0fBDF8CDCC; | |
| fma.rn.ftz.f32 %f107, %f105, %f102, %f106; | |
| mov.f32 %f108, 0f3E0F2955; | |
| fma.rn.ftz.f32 %f109, %f107, %f102, %f108; | |
| mov.f32 %f110, 0fBE2AD8B9; | |
| fma.rn.ftz.f32 %f111, %f109, %f102, %f110; | |
| mov.f32 %f112, 0f3E4CED0B; | |
| fma.rn.ftz.f32 %f113, %f111, %f102, %f112; | |
| mov.f32 %f114, 0fBE7FFF22; | |
| fma.rn.ftz.f32 %f115, %f113, %f102, %f114; | |
| mov.f32 %f116, 0f3EAAAA78; | |
| fma.rn.ftz.f32 %f117, %f115, %f102, %f116; | |
| mov.f32 %f118, 0fBF000000; | |
| fma.rn.ftz.f32 %f119, %f117, %f102, %f118; | |
| mul.f32 %f120, %f102, %f119; | |
| fma.rn.ftz.f32 %f121, %f120, %f102, %f102; | |
| mov.f32 %f122, 0f3F317218; | |
| fma.rn.ftz.f32 %f393, %f101, %f122, %f121; | |
| setp.lt.u32 %p102, %r140, 2139095040; | |
| mov.f32 %f123, 0f7F800000; | |
| @%p102 bra $L__BB0_5; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f393, %f33, %f123, %f123; | |
| $L__BB0_5: | |
| mov.b32 %f18, %r83; | |
| .loc 1 50 23 | |
| setp.lt.f32 %p104, %f18, 0f00800000; | |
| mul.f32 %f124, %f18, 0f4B000000; | |
| selp.f32 %f38, %f124, %f18, %p104; | |
| selp.f32 %f125, 0fC1B80000, 0f00000000, %p104; | |
| mov.b32 %r144, %f38; | |
| add.s32 %r145, %r144, -1059760811; | |
| and.b32 %r146, %r145, -8388608; | |
| sub.s32 %r147, %r144, %r146; | |
| mov.b32 %f126, %r147; | |
| cvt.rn.f32.s32 %f127, %r146; | |
| fma.rn.ftz.f32 %f129, %f127, %f100, %f125; | |
| add.f32 %f130, %f126, 0fBF800000; | |
| fma.rn.ftz.f32 %f133, %f104, %f130, %f103; | |
| fma.rn.ftz.f32 %f135, %f133, %f130, %f106; | |
| fma.rn.ftz.f32 %f137, %f135, %f130, %f108; | |
| fma.rn.ftz.f32 %f139, %f137, %f130, %f110; | |
| fma.rn.ftz.f32 %f141, %f139, %f130, %f112; | |
| fma.rn.ftz.f32 %f143, %f141, %f130, %f114; | |
| fma.rn.ftz.f32 %f145, %f143, %f130, %f116; | |
| fma.rn.ftz.f32 %f147, %f145, %f130, %f118; | |
| mul.f32 %f148, %f130, %f147; | |
| fma.rn.ftz.f32 %f149, %f148, %f130, %f130; | |
| fma.rn.ftz.f32 %f394, %f129, %f122, %f149; | |
| setp.lt.u32 %p105, %r144, 2139095040; | |
| @%p105 bra $L__BB0_7; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f394, %f38, %f123, %f123; | |
| $L__BB0_7: | |
| mov.b32 %f19, %r84; | |
| .loc 1 50 23 | |
| setp.lt.f32 %p107, %f19, 0f00800000; | |
| mul.f32 %f152, %f19, 0f4B000000; | |
| selp.f32 %f43, %f152, %f19, %p107; | |
| selp.f32 %f153, 0fC1B80000, 0f00000000, %p107; | |
| mov.b32 %r148, %f43; | |
| add.s32 %r149, %r148, -1059760811; | |
| and.b32 %r150, %r149, -8388608; | |
| sub.s32 %r151, %r148, %r150; | |
| mov.b32 %f154, %r151; | |
| cvt.rn.f32.s32 %f155, %r150; | |
| fma.rn.ftz.f32 %f157, %f155, %f100, %f153; | |
| add.f32 %f158, %f154, 0fBF800000; | |
| fma.rn.ftz.f32 %f161, %f104, %f158, %f103; | |
| fma.rn.ftz.f32 %f163, %f161, %f158, %f106; | |
| fma.rn.ftz.f32 %f165, %f163, %f158, %f108; | |
| fma.rn.ftz.f32 %f167, %f165, %f158, %f110; | |
| fma.rn.ftz.f32 %f169, %f167, %f158, %f112; | |
| fma.rn.ftz.f32 %f171, %f169, %f158, %f114; | |
| fma.rn.ftz.f32 %f173, %f171, %f158, %f116; | |
| fma.rn.ftz.f32 %f175, %f173, %f158, %f118; | |
| mul.f32 %f176, %f158, %f175; | |
| fma.rn.ftz.f32 %f177, %f176, %f158, %f158; | |
| fma.rn.ftz.f32 %f395, %f157, %f122, %f177; | |
| setp.lt.u32 %p108, %r148, 2139095040; | |
| @%p108 bra $L__BB0_9; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f395, %f43, %f123, %f123; | |
| $L__BB0_9: | |
| mov.b32 %f20, %r85; | |
| .loc 1 50 23 | |
| setp.lt.f32 %p110, %f20, 0f00800000; | |
| mul.f32 %f180, %f20, 0f4B000000; | |
| selp.f32 %f48, %f180, %f20, %p110; | |
| selp.f32 %f181, 0fC1B80000, 0f00000000, %p110; | |
| mov.b32 %r152, %f48; | |
| add.s32 %r153, %r152, -1059760811; | |
| and.b32 %r154, %r153, -8388608; | |
| sub.s32 %r155, %r152, %r154; | |
| mov.b32 %f182, %r155; | |
| cvt.rn.f32.s32 %f183, %r154; | |
| fma.rn.ftz.f32 %f185, %f183, %f100, %f181; | |
| add.f32 %f186, %f182, 0fBF800000; | |
| fma.rn.ftz.f32 %f189, %f104, %f186, %f103; | |
| fma.rn.ftz.f32 %f191, %f189, %f186, %f106; | |
| fma.rn.ftz.f32 %f193, %f191, %f186, %f108; | |
| fma.rn.ftz.f32 %f195, %f193, %f186, %f110; | |
| fma.rn.ftz.f32 %f197, %f195, %f186, %f112; | |
| fma.rn.ftz.f32 %f199, %f197, %f186, %f114; | |
| fma.rn.ftz.f32 %f201, %f199, %f186, %f116; | |
| fma.rn.ftz.f32 %f203, %f201, %f186, %f118; | |
| mul.f32 %f204, %f186, %f203; | |
| fma.rn.ftz.f32 %f205, %f204, %f186, %f186; | |
| fma.rn.ftz.f32 %f396, %f185, %f122, %f205; | |
| setp.lt.u32 %p111, %r152, 2139095040; | |
| @%p111 bra $L__BB0_11; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f396, %f48, %f123, %f123; | |
| $L__BB0_11: | |
| mov.b32 %f21, %r90; | |
| .loc 1 50 23 | |
| setp.lt.f32 %p113, %f21, 0f00800000; | |
| mul.f32 %f208, %f21, 0f4B000000; | |
| selp.f32 %f53, %f208, %f21, %p113; | |
| selp.f32 %f209, 0fC1B80000, 0f00000000, %p113; | |
| mov.b32 %r156, %f53; | |
| add.s32 %r157, %r156, -1059760811; | |
| and.b32 %r158, %r157, -8388608; | |
| sub.s32 %r159, %r156, %r158; | |
| mov.b32 %f210, %r159; | |
| cvt.rn.f32.s32 %f211, %r158; | |
| fma.rn.ftz.f32 %f213, %f211, %f100, %f209; | |
| add.f32 %f214, %f210, 0fBF800000; | |
| fma.rn.ftz.f32 %f217, %f104, %f214, %f103; | |
| fma.rn.ftz.f32 %f219, %f217, %f214, %f106; | |
| fma.rn.ftz.f32 %f221, %f219, %f214, %f108; | |
| fma.rn.ftz.f32 %f223, %f221, %f214, %f110; | |
| fma.rn.ftz.f32 %f225, %f223, %f214, %f112; | |
| fma.rn.ftz.f32 %f227, %f225, %f214, %f114; | |
| fma.rn.ftz.f32 %f229, %f227, %f214, %f116; | |
| fma.rn.ftz.f32 %f231, %f229, %f214, %f118; | |
| mul.f32 %f232, %f214, %f231; | |
| fma.rn.ftz.f32 %f233, %f232, %f214, %f214; | |
| fma.rn.ftz.f32 %f397, %f213, %f122, %f233; | |
| setp.lt.u32 %p114, %r156, 2139095040; | |
| @%p114 bra $L__BB0_13; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f397, %f53, %f123, %f123; | |
| $L__BB0_13: | |
| mov.b32 %f22, %r91; | |
| .loc 1 50 23 | |
| setp.lt.f32 %p116, %f22, 0f00800000; | |
| mul.f32 %f236, %f22, 0f4B000000; | |
| selp.f32 %f58, %f236, %f22, %p116; | |
| selp.f32 %f237, 0fC1B80000, 0f00000000, %p116; | |
| mov.b32 %r160, %f58; | |
| add.s32 %r161, %r160, -1059760811; | |
| and.b32 %r162, %r161, -8388608; | |
| sub.s32 %r163, %r160, %r162; | |
| mov.b32 %f238, %r163; | |
| cvt.rn.f32.s32 %f239, %r162; | |
| fma.rn.ftz.f32 %f241, %f239, %f100, %f237; | |
| add.f32 %f242, %f238, 0fBF800000; | |
| fma.rn.ftz.f32 %f245, %f104, %f242, %f103; | |
| fma.rn.ftz.f32 %f247, %f245, %f242, %f106; | |
| fma.rn.ftz.f32 %f249, %f247, %f242, %f108; | |
| fma.rn.ftz.f32 %f251, %f249, %f242, %f110; | |
| fma.rn.ftz.f32 %f253, %f251, %f242, %f112; | |
| fma.rn.ftz.f32 %f255, %f253, %f242, %f114; | |
| fma.rn.ftz.f32 %f257, %f255, %f242, %f116; | |
| fma.rn.ftz.f32 %f259, %f257, %f242, %f118; | |
| mul.f32 %f260, %f242, %f259; | |
| fma.rn.ftz.f32 %f261, %f260, %f242, %f242; | |
| fma.rn.ftz.f32 %f398, %f241, %f122, %f261; | |
| setp.lt.u32 %p117, %r160, 2139095040; | |
| @%p117 bra $L__BB0_15; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f398, %f58, %f123, %f123; | |
| $L__BB0_15: | |
| and.pred %p52, %p16, %p46; | |
| and.pred %p53, %p16, %p47; | |
| and.pred %p54, %p16, %p44; | |
| and.pred %p55, %p16, %p45; | |
| and.pred %p56, %p1, %p50; | |
| and.pred %p57, %p1, %p51; | |
| and.pred %p58, %p1, %p48; | |
| and.pred %p59, %p1, %p49; | |
| setp.eq.f32 %p103, %f33, 0f00000000; | |
| setp.eq.f32 %p106, %f38, 0f00000000; | |
| setp.eq.f32 %p109, %f43, 0f00000000; | |
| setp.eq.f32 %p112, %f48, 0f00000000; | |
| setp.eq.f32 %p115, %f53, 0f00000000; | |
| mov.b32 %f23, %r92; | |
| .loc 1 50 23 | |
| setp.eq.f32 %p118, %f58, 0f00000000; | |
| setp.lt.f32 %p119, %f23, 0f00800000; | |
| mul.f32 %f264, %f23, 0f4B000000; | |
| selp.f32 %f63, %f264, %f23, %p119; | |
| selp.f32 %f265, 0fC1B80000, 0f00000000, %p119; | |
| mov.b32 %r164, %f63; | |
| add.s32 %r165, %r164, -1059760811; | |
| and.b32 %r166, %r165, -8388608; | |
| sub.s32 %r167, %r164, %r166; | |
| mov.b32 %f266, %r167; | |
| cvt.rn.f32.s32 %f267, %r166; | |
| fma.rn.ftz.f32 %f269, %f267, %f100, %f265; | |
| add.f32 %f270, %f266, 0fBF800000; | |
| fma.rn.ftz.f32 %f273, %f104, %f270, %f103; | |
| fma.rn.ftz.f32 %f275, %f273, %f270, %f106; | |
| fma.rn.ftz.f32 %f277, %f275, %f270, %f108; | |
| fma.rn.ftz.f32 %f279, %f277, %f270, %f110; | |
| fma.rn.ftz.f32 %f281, %f279, %f270, %f112; | |
| fma.rn.ftz.f32 %f283, %f281, %f270, %f114; | |
| fma.rn.ftz.f32 %f285, %f283, %f270, %f116; | |
| fma.rn.ftz.f32 %f287, %f285, %f270, %f118; | |
| mul.f32 %f288, %f270, %f287; | |
| fma.rn.ftz.f32 %f289, %f288, %f270, %f270; | |
| fma.rn.ftz.f32 %f399, %f269, %f122, %f289; | |
| setp.lt.u32 %p120, %r164, 2139095040; | |
| @%p120 bra $L__BB0_17; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f399, %f63, %f123, %f123; | |
| $L__BB0_17: | |
| selp.u64 %rd78, 1, 0, %p55; | |
| selp.u64 %rd79, 1, 0, %p54; | |
| selp.u64 %rd80, 1, 0, %p53; | |
| selp.u64 %rd81, 1, 0, %p52; | |
| selp.u64 %rd82, 1, 0, %p59; | |
| selp.u64 %rd83, 1, 0, %p58; | |
| selp.u64 %rd84, 1, 0, %p57; | |
| selp.u64 %rd85, 1, 0, %p56; | |
| mov.b32 %f24, %r93; | |
| cvt.f32.bf16 %r131, %rs43; | |
| mov.b32 %f88, %r131; | |
| cvt.f32.bf16 %r132, %rs45; | |
| mov.b32 %f89, %r132; | |
| cvt.f32.bf16 %r133, %rs47; | |
| mov.b32 %f90, %r133; | |
| cvt.f32.bf16 %r134, %rs49; | |
| mov.b32 %f91, %r134; | |
| cvt.f32.bf16 %r135, %rs51; | |
| mov.b32 %f92, %r135; | |
| cvt.f32.bf16 %r136, %rs53; | |
| mov.b32 %f93, %r136; | |
| cvt.f32.bf16 %r137, %rs55; | |
| mov.b32 %f94, %r137; | |
| cvt.f32.bf16 %r138, %rs57; | |
| mov.b32 %f95, %r138; | |
| sub.f32 %f32, %f95, %f16; | |
| sub.f32 %f31, %f94, %f15; | |
| sub.f32 %f30, %f93, %f14; | |
| sub.f32 %f29, %f92, %f13; | |
| sub.f32 %f28, %f91, %f12; | |
| sub.f32 %f27, %f90, %f11; | |
| sub.f32 %f26, %f89, %f10; | |
| sub.f32 %f25, %f88, %f9; | |
| .loc 1 50 23 | |
| selp.f32 %f37, 0fFF800000, %f393, %p103; | |
| selp.f32 %f42, 0fFF800000, %f394, %p106; | |
| selp.f32 %f47, 0fFF800000, %f395, %p109; | |
| selp.f32 %f52, 0fFF800000, %f396, %p112; | |
| selp.f32 %f57, 0fFF800000, %f397, %p115; | |
| selp.f32 %f62, 0fFF800000, %f398, %p118; | |
| setp.eq.f32 %p121, %f63, 0f00000000; | |
| selp.f32 %f67, 0fFF800000, %f399, %p121; | |
| setp.lt.f32 %p122, %f24, 0f00800000; | |
| mul.f32 %f292, %f24, 0f4B000000; | |
| selp.f32 %f68, %f292, %f24, %p122; | |
| selp.f32 %f293, 0fC1B80000, 0f00000000, %p122; | |
| mov.b32 %r168, %f68; | |
| add.s32 %r169, %r168, -1059760811; | |
| and.b32 %r170, %r169, -8388608; | |
| sub.s32 %r171, %r168, %r170; | |
| mov.b32 %f294, %r171; | |
| cvt.rn.f32.s32 %f295, %r170; | |
| fma.rn.ftz.f32 %f297, %f295, %f100, %f293; | |
| add.f32 %f298, %f294, 0fBF800000; | |
| fma.rn.ftz.f32 %f301, %f104, %f298, %f103; | |
| fma.rn.ftz.f32 %f303, %f301, %f298, %f106; | |
| fma.rn.ftz.f32 %f305, %f303, %f298, %f108; | |
| fma.rn.ftz.f32 %f307, %f305, %f298, %f110; | |
| fma.rn.ftz.f32 %f309, %f307, %f298, %f112; | |
| fma.rn.ftz.f32 %f311, %f309, %f298, %f114; | |
| fma.rn.ftz.f32 %f313, %f311, %f298, %f116; | |
| fma.rn.ftz.f32 %f315, %f313, %f298, %f118; | |
| mul.f32 %f316, %f298, %f315; | |
| fma.rn.ftz.f32 %f317, %f316, %f298, %f298; | |
| fma.rn.ftz.f32 %f400, %f297, %f122, %f317; | |
| setp.lt.u32 %p123, %r168, 2139095040; | |
| @%p123 bra $L__BB0_19; | |
| .loc 1 0 23 | |
| fma.rn.ftz.f32 %f400, %f68, %f123, %f123; | |
| bra.uni $L__BB0_19; | |
| $L__BB0_20: | |
| .loc 1 24 33 | |
| bfe.u32 %r191, %r1, 5, 3; | |
| and.b32 %r192, %r1, 31; | |
| $L__tmp1: | |
| .loc 2 243 36 | |
| bar.sync 0; | |
| $L__tmp2: | |
| .loc 2 233 15 | |
| add.s64 %rd146, %rd201, %rd202; | |
| add.s64 %rd147, %rd146, %rd203; | |
| add.s64 %rd148, %rd147, %rd204; | |
| add.s64 %rd149, %rd148, %rd205; | |
| add.s64 %rd150, %rd149, %rd206; | |
| add.s64 %rd151, %rd150, %rd207; | |
| add.s64 %rd152, %rd151, %rd208; | |
| $L__tmp3: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r193, %rd152; | |
| shfl.sync.bfly.b32 %r194, %r193, 16, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r195}, %rd152; } | |
| shfl.sync.bfly.b32 %r196, %r195, 16, 31, -1; | |
| cvt.u64.u32 %rd153, %r194; | |
| cvt.u64.u32 %rd154, %r196; | |
| shl.b64 %rd155, %rd154, 32; | |
| or.b64 %rd156, %rd153, %rd155; | |
| $L__tmp4: | |
| .loc 2 233 15 | |
| add.s64 %rd157, %rd152, %rd156; | |
| $L__tmp5: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r197, %rd157; | |
| shfl.sync.bfly.b32 %r198, %r197, 8, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r199}, %rd157; } | |
| shfl.sync.bfly.b32 %r200, %r199, 8, 31, -1; | |
| cvt.u64.u32 %rd158, %r198; | |
| cvt.u64.u32 %rd159, %r200; | |
| shl.b64 %rd160, %rd159, 32; | |
| or.b64 %rd161, %rd158, %rd160; | |
| $L__tmp6: | |
| .loc 2 233 15 | |
| add.s64 %rd162, %rd157, %rd161; | |
| $L__tmp7: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r201, %rd162; | |
| shfl.sync.bfly.b32 %r202, %r201, 4, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r203}, %rd162; } | |
| shfl.sync.bfly.b32 %r204, %r203, 4, 31, -1; | |
| cvt.u64.u32 %rd163, %r202; | |
| cvt.u64.u32 %rd164, %r204; | |
| shl.b64 %rd165, %rd164, 32; | |
| or.b64 %rd166, %rd163, %rd165; | |
| $L__tmp8: | |
| .loc 2 233 15 | |
| add.s64 %rd167, %rd162, %rd166; | |
| $L__tmp9: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r205, %rd167; | |
| shfl.sync.bfly.b32 %r206, %r205, 2, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r207}, %rd167; } | |
| shfl.sync.bfly.b32 %r208, %r207, 2, 31, -1; | |
| cvt.u64.u32 %rd168, %r206; | |
| cvt.u64.u32 %rd169, %r208; | |
| shl.b64 %rd170, %rd169, 32; | |
| or.b64 %rd171, %rd168, %rd170; | |
| $L__tmp10: | |
| .loc 2 233 15 | |
| add.s64 %rd172, %rd167, %rd171; | |
| $L__tmp11: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r209, %rd172; | |
| shfl.sync.bfly.b32 %r210, %r209, 1, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r211}, %rd172; } | |
| shfl.sync.bfly.b32 %r212, %r211, 1, 31, -1; | |
| cvt.u64.u32 %rd173, %r210; | |
| cvt.u64.u32 %rd174, %r212; | |
| shl.b64 %rd175, %rd174, 32; | |
| or.b64 %rd176, %rd173, %rd175; | |
| $L__tmp12: | |
| .loc 2 233 15 | |
| add.s64 %rd140, %rd172, %rd176; | |
| $L__tmp13: | |
| .loc 2 243 36 | |
| setp.eq.s32 %p134, %r192, 0; | |
| shl.b32 %r213, %r191, 3; | |
| add.s32 %r180, %r44, %r213; | |
| @%p134 st.shared.b64 [ %r180 + 0 ], %rd140; | |
| bar.sync 0; | |
| setp.lt.s32 %p135, %r1, 8; | |
| shl.b32 %r215, %r1, 3; | |
| add.s32 %r181, %r44, %r215; | |
| @%p135 ld.shared.b64 %rd141, [ %r181 + 0 ]; | |
| cvt.u32.u64 %r216, %rd141; | |
| shfl.sync.bfly.b32 %r217, %r216, 4, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r218}, %rd141; } | |
| shfl.sync.bfly.b32 %r219, %r218, 4, 31, -1; | |
| cvt.u64.u32 %rd177, %r217; | |
| cvt.u64.u32 %rd178, %r219; | |
| shl.b64 %rd179, %rd178, 32; | |
| or.b64 %rd180, %rd177, %rd179; | |
| $L__tmp14: | |
| .loc 2 233 15 | |
| add.s64 %rd181, %rd141, %rd180; | |
| $L__tmp15: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r220, %rd181; | |
| shfl.sync.bfly.b32 %r221, %r220, 2, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r222}, %rd181; } | |
| shfl.sync.bfly.b32 %r223, %r222, 2, 31, -1; | |
| cvt.u64.u32 %rd182, %r221; | |
| cvt.u64.u32 %rd183, %r223; | |
| shl.b64 %rd184, %rd183, 32; | |
| or.b64 %rd185, %rd182, %rd184; | |
| $L__tmp16: | |
| .loc 2 233 15 | |
| add.s64 %rd186, %rd181, %rd185; | |
| $L__tmp17: | |
| .loc 2 243 36 | |
| cvt.u32.u64 %r224, %rd186; | |
| shfl.sync.bfly.b32 %r225, %r224, 1, 31, -1; | |
| { .reg .b32 tmp; mov.b64 {tmp, %r226}, %rd186; } | |
| shfl.sync.bfly.b32 %r227, %r226, 1, 31, -1; | |
| cvt.u64.u32 %rd187, %r225; | |
| cvt.u64.u32 %rd188, %r227; | |
| shl.b64 %rd189, %rd188, 32; | |
| or.b64 %rd190, %rd187, %rd189; | |
| $L__tmp18: | |
| .loc 2 233 15 | |
| add.s64 %rd142, %rd186, %rd190; | |
| $L__tmp19: | |
| .loc 2 243 36 | |
| and.b32 %r228, %r1, 7; | |
| setp.eq.s32 %p143, %r228, 0; | |
| and.pred %p136, %p135, %p143; | |
| @%p136 st.shared.b64 [ %r181 + 0 ], %rd142; | |
| bar.sync 0; | |
| ld.shared.u32 %rd191, [global_smem+4]; | |
| shl.b64 %rd192, %rd191, 32; | |
| ld.shared.u32 %rd193, [global_smem]; | |
| or.b64 %rd194, %rd192, %rd193; | |
| $L__tmp20: | |
| .loc 1 60 28 | |
| bar.sync 0; | |
| st.shared.u64 [global_smem], %rd194; | |
| bar.sync 0; | |
| ld.shared.u64 %rd143, [global_smem]; | |
| .loc 1 61 25 | |
| shl.b64 %rd195, %rd1, 3; | |
| add.s64 %rd144, %rd39, %rd195; | |
| .loc 1 61 36 | |
| setp.eq.s32 %p144, %r2, 0; | |
| and.pred %p137, %p144, %p85; | |
| @%p137 st.global.b64 [ %rd144 + 0 ], { %rd143 }; | |
| $L__tmp21: | |
| .loc 2 243 36 | |
| bar.sync 0; | |
| $L__tmp22: | |
| .loc 2 233 15 | |
| add.f32 %f361, %f385, %f386; | |
| add.f32 %f362, %f387, %f361; | |
| add.f32 %f363, %f388, %f362; | |
| add.f32 %f364, %f389, %f363; | |
| add.f32 %f365, %f390, %f364; | |
| add.f32 %f366, %f391, %f365; | |
| add.f32 %f367, %f392, %f366; | |
| $L__tmp23: | |
| .loc 2 243 36 | |
| mov.b32 %r229, %f367; | |
| shfl.sync.bfly.b32 %r230, %r229, 16, 31, -1; | |
| mov.b32 %f368, %r230; | |
| $L__tmp24: | |
| .loc 2 233 15 | |
| add.f32 %f369, %f367, %f368; | |
| $L__tmp25: | |
| .loc 2 243 36 | |
| mov.b32 %r231, %f369; | |
| shfl.sync.bfly.b32 %r232, %r231, 8, 31, -1; | |
| mov.b32 %f370, %r232; | |
| $L__tmp26: | |
| .loc 2 233 15 | |
| add.f32 %f371, %f369, %f370; | |
| $L__tmp27: | |
| .loc 2 243 36 | |
| mov.b32 %r233, %f371; | |
| shfl.sync.bfly.b32 %r234, %r233, 4, 31, -1; | |
| mov.b32 %f372, %r234; | |
| $L__tmp28: | |
| .loc 2 233 15 | |
| add.f32 %f373, %f371, %f372; | |
| $L__tmp29: | |
| .loc 2 243 36 | |
| mov.b32 %r235, %f373; | |
| shfl.sync.bfly.b32 %r236, %r235, 2, 31, -1; | |
| mov.b32 %f374, %r236; | |
| $L__tmp30: | |
| .loc 2 233 15 | |
| add.f32 %f375, %f373, %f374; | |
| $L__tmp31: | |
| .loc 2 243 36 | |
| mov.b32 %r237, %f375; | |
| shfl.sync.bfly.b32 %r238, %r237, 1, 31, -1; | |
| mov.b32 %f376, %r238; | |
| $L__tmp32: | |
| .loc 2 233 15 | |
| add.f32 %f377, %f375, %f376; | |
| $L__tmp33: | |
| .loc 2 243 36 | |
| shl.b32 %r239, %r191, 2; | |
| add.s32 %r183, %r44, %r239; | |
| mov.b32 %r184, %f377; | |
| @%p134 st.shared.b32 [ %r183 + 0 ], %r184; | |
| bar.sync 0; | |
| shl.b32 %r240, %r1, 2; | |
| add.s32 %r186, %r44, %r240; | |
| @%p135 ld.shared.b32 %r185, [ %r186 + 0 ]; | |
| mov.b32 %f378, %r185; | |
| shfl.sync.bfly.b32 %r241, %r185, 4, 31, -1; | |
| mov.b32 %f379, %r241; | |
| $L__tmp34: | |
| .loc 2 233 15 | |
| add.f32 %f380, %f378, %f379; | |
| $L__tmp35: | |
| .loc 2 243 36 | |
| mov.b32 %r242, %f380; | |
| shfl.sync.bfly.b32 %r243, %r242, 2, 31, -1; | |
| mov.b32 %f381, %r243; | |
| $L__tmp36: | |
| .loc 2 233 15 | |
| add.f32 %f382, %f380, %f381; | |
| $L__tmp37: | |
| .loc 2 243 36 | |
| mov.b32 %r244, %f382; | |
| shfl.sync.bfly.b32 %r245, %r244, 1, 31, -1; | |
| mov.b32 %f383, %r245; | |
| $L__tmp38: | |
| .loc 2 233 15 | |
| add.f32 %f384, %f382, %f383; | |
| $L__tmp39: | |
| .loc 2 243 36 | |
| mov.b32 %r188, %f384; | |
| @%p136 st.shared.b32 [ %r186 + 0 ], %r188; | |
| bar.sync 0; | |
| ld.shared.u32 %r189, [global_smem]; | |
| $L__tmp40: | |
| .loc 1 63 25 | |
| shl.b64 %rd196, %rd1, 2; | |
| add.s64 %rd145, %rd40, %rd196; | |
| .loc 1 63 37 | |
| @%p137 st.global.b32 [ %rd145 + 0 ], { %r189 }; | |
| .loc 1 63 4 | |
| ret; | |
| $L__tmp41: | |
| $L__func_end0: | |
| } | |
| // .globl __nv_logf | |
| .visible .func (.param .b32 func_retval0) __nv_logf( | |
| .param .b32 __nv_logf_param_0 | |
| ) | |
| { | |
| .reg .pred %p<4>; | |
| .reg .b32 %r<5>; | |
| .reg .f32 %f<36>; | |
| $L__func_begin1: | |
| ld.param.f32 %f5, [__nv_logf_param_0]; | |
| setp.lt.f32 %p1, %f5, 0f00800000; | |
| mul.f32 %f6, %f5, 0f4B000000; | |
| selp.f32 %f1, %f6, %f5, %p1; | |
| selp.f32 %f7, 0fC1B80000, 0f00000000, %p1; | |
| mov.b32 %r1, %f1; | |
| add.s32 %r2, %r1, -1059760811; | |
| and.b32 %r3, %r2, -8388608; | |
| sub.s32 %r4, %r1, %r3; | |
| mov.b32 %f8, %r4; | |
| cvt.rn.f32.s32 %f9, %r3; | |
| mov.f32 %f10, 0f34000000; | |
| fma.rn.ftz.f32 %f11, %f9, %f10, %f7; | |
| add.f32 %f12, %f8, 0fBF800000; | |
| mov.f32 %f13, 0f3E1039F6; | |
| mov.f32 %f14, 0fBE055027; | |
| fma.rn.ftz.f32 %f15, %f14, %f12, %f13; | |
| mov.f32 %f16, 0fBDF8CDCC; | |
| fma.rn.ftz.f32 %f17, %f15, %f12, %f16; | |
| mov.f32 %f18, 0f3E0F2955; | |
| fma.rn.ftz.f32 %f19, %f17, %f12, %f18; | |
| mov.f32 %f20, 0fBE2AD8B9; | |
| fma.rn.ftz.f32 %f21, %f19, %f12, %f20; | |
| mov.f32 %f22, 0f3E4CED0B; | |
| fma.rn.ftz.f32 %f23, %f21, %f12, %f22; | |
| mov.f32 %f24, 0fBE7FFF22; | |
| fma.rn.ftz.f32 %f25, %f23, %f12, %f24; | |
| mov.f32 %f26, 0f3EAAAA78; | |
| fma.rn.ftz.f32 %f27, %f25, %f12, %f26; | |
| mov.f32 %f28, 0fBF000000; | |
| fma.rn.ftz.f32 %f29, %f27, %f12, %f28; | |
| mul.f32 %f30, %f12, %f29; | |
| fma.rn.ftz.f32 %f31, %f30, %f12, %f12; | |
| mov.f32 %f32, 0f3F317218; | |
| fma.rn.ftz.f32 %f35, %f11, %f32, %f31; | |
| setp.lt.u32 %p2, %r1, 2139095040; | |
| @%p2 bra $L__BB1_2; | |
| mov.f32 %f33, 0f7F800000; | |
| fma.rn.ftz.f32 %f35, %f1, %f33, %f33; | |
| $L__BB1_2: | |
| setp.eq.f32 %p3, %f1, 0f00000000; | |
| selp.f32 %f34, 0fFF800000, %f35, %p3; | |
| st.param.f32 [func_retval0+0], %f34; | |
| ret; | |
| $L__func_end1: | |
| } | |
| .file 1 "/tmp/torchinductor_root/ul/culwqy52mqs4o2bmqocf2r5plomw2phviv5gutbxlcpdrdkc46ri.py" | |
| .file 2 "/usr/local/lib/python3.10/dist-packages/triton/language/standard.py" | |
| .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 180 | |
| .b8 66 | |
| .b8 12 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 0 | |
| .b8 0 | |
| .b8 2 | |
| .b8 46 | |
| .b8 0 | |
| .b8 135 | |
| .b8 64 | |
| .b8 8 | |
| .b8 3 | |
| .b8 8 | |
| .b8 58 | |
| .b8 11 | |
| .b8 59 | |
| .b8 11 | |
| .b8 63 | |
| .b8 12 | |
| .b8 32 | |
| .b8 11 | |
| .b8 0 | |
| .b8 0 | |
| .b8 3 | |
| .b8 46 | |
| .b8 1 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 64 | |
| .b8 10 | |
| .b8 49 | |
| .b8 19 | |
| .b8 0 | |
| .b8 0 | |
| .b8 4 | |
| .b8 29 | |
| .b8 0 | |
| .b8 49 | |
| .b8 19 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 88 | |
| .b8 11 | |
| .b8 89 | |
| .b8 11 | |
| .b8 87 | |
| .b8 11 | |
| .b8 0 | |
| .b8 0 | |
| .b8 5 | |
| .b8 29 | |
| .b8 1 | |
| .b8 49 | |
| .b8 19 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 88 | |
| .b8 11 | |
| .b8 89 | |
| .b8 11 | |
| .b8 87 | |
| .b8 11 | |
| .b8 0 | |
| .b8 0 | |
| .b8 0 | |
| } | |
| .section .debug_info | |
| { | |
| .b32 349 | |
| .b8 2 | |
| .b8 0 | |
| .b32 .debug_abbrev | |
| .b8 8 | |
| .b8 1 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 0 | |
| .b8 2 | |
| .b8 0 | |
| .b8 99 | |
| .b8 117 | |
| .b8 108 | |
| .b8 119 | |
| .b8 113 | |
| .b8 121 | |
| .b8 53 | |
| .b8 50 | |
| .b8 109 | |
| .b8 113 | |
| .b8 115 | |
| .b8 52 | |
| .b8 111 | |
| .b8 50 | |
| .b8 98 | |
| .b8 109 | |
| .b8 113 | |
| .b8 111 | |
| .b8 99 | |
| .b8 102 | |
| .b8 50 | |
| .b8 114 | |
| .b8 53 | |
| .b8 112 | |
| .b8 108 | |
| .b8 111 | |
| .b8 109 | |
| .b8 119 | |
| .b8 50 | |
| .b8 112 | |
| .b8 104 | |
| .b8 118 | |
| .b8 105 | |
| .b8 118 | |
| .b8 53 | |
| .b8 103 | |
| .b8 117 | |
| .b8 116 | |
| .b8 98 | |
| .b8 120 | |
| .b8 108 | |
| .b8 99 | |
| .b8 112 | |
| .b8 100 | |
| .b8 114 | |
| .b8 100 | |
| .b8 107 | |
| .b8 99 | |
| .b8 52 | |
| .b8 54 | |
| .b8 114 | |
| .b8 105 | |
| .b8 46 | |
| .b8 112 | |
| .b8 121 | |
| .b8 0 | |
| .b32 .debug_line | |
| .b8 47 | |
| .b8 116 | |
| .b8 109 | |
| .b8 112 | |
| .b8 47 | |
| .b8 116 | |
| .b8 111 | |
| .b8 114 | |
| .b8 99 | |
| .b8 104 | |
| .b8 105 | |
| .b8 110 | |
| .b8 100 | |
| .b8 117 | |
| .b8 99 | |
| .b8 116 | |
| .b8 111 | |
| .b8 114 | |
| .b8 95 | |
| .b8 114 | |
| .b8 111 | |
| .b8 111 | |
| .b8 116 | |
| .b8 47 | |
| .b8 117 | |
| .b8 108 | |
| .b8 0 | |
| .b8 1 | |
| .b64 $L__func_begin0 | |
| .b64 $L__func_end0 | |
| .b8 2 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 95 | |
| .b8 95 | |
| .b8 48 | |
| .b8 100 | |
| .b8 49 | |
| .b8 100 | |
| .b8 50 | |
| .b8 100 | |
| .b8 51 | |
| .b8 100 | |
| .b8 52 | |
| .b8 100 | |
| .b8 53 | |
| .b8 100 | |
| .b8 54 | |
| .b8 101 | |
| .b8 55 | |
| .b8 100 | |
| .b8 101 | |
| .b8 0 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 95 | |
| .b8 95 | |
| .b8 48 | |
| .b8 100 | |
| .b8 49 | |
| .b8 100 | |
| .b8 50 | |
| .b8 100 | |
| .b8 51 | |
| .b8 100 | |
| .b8 52 | |
| .b8 100 | |
| .b8 53 | |
| .b8 100 | |
| .b8 54 | |
| .b8 101 | |
| .b8 55 | |
| .b8 100 | |
| .b8 101 | |
| .b8 0 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 1 | |
| .b8 3 | |
| .b64 $L__func_begin0 | |
| .b64 $L__func_end0 | |
| .b8 1 | |
| .b8 156 | |
| .b32 125 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp1 | |
| .b64 $L__tmp20 | |
| .b8 2 | |
| .b8 60 | |
| .b8 25 | |
| .b8 5 | |
| .b32 125 | |
| .b64 $L__tmp2 | |
| .b64 $L__tmp19 | |
| .b8 2 | |
| .b8 60 | |
| .b8 25 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp2 | |
| .b64 $L__tmp19 | |
| .b8 2 | |
| .b8 243 | |
| .b8 36 | |
| .b8 0 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp21 | |
| .b64 $L__tmp40 | |
| .b8 2 | |
| .b8 62 | |
| .b8 27 | |
| .b8 5 | |
| .b32 125 | |
| .b64 $L__tmp22 | |
| .b64 $L__tmp39 | |
| .b8 2 | |
| .b8 62 | |
| .b8 27 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp22 | |
| .b64 $L__tmp39 | |
| .b8 2 | |
| .b8 243 | |
| .b8 36 | |
| .b8 0 | |
| .b8 0 | |
| .b8 0 | |
| } | |
| .section .debug_pubnames | |
| { | |
| .b32 $L__pubNames_end0-$L__pubNames_start0 | |
| $L__pubNames_start0: | |
| .b8 2 | |
| .b8 0 | |
| .b32 .debug_info | |
| .b32 353 | |
| .b32 125 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 95 | |
| .b8 95 | |
| .b8 48 | |
| .b8 100 | |
| .b8 49 | |
| .b8 100 | |
| .b8 50 | |
| .b8 100 | |
| .b8 51 | |
| .b8 100 | |
| .b8 52 | |
| .b8 100 | |
| .b8 53 | |
| .b8 100 | |
| .b8 54 | |
| .b8 101 | |
| .b8 55 | |
| .b8 100 | |
| .b8 101 | |
| .b8 0 | |
| .b32 0 | |
| $L__pubNames_end0: | |
| } | |
| .section .debug_pubtypes | |
| { | |
| .b32 $L__pubTypes_end0-$L__pubTypes_start0 | |
| $L__pubTypes_start0: | |
| .b8 2 | |
| .b8 0 | |
| .b32 .debug_info | |
| .b32 353 | |
| .b32 0 | |
| $L__pubTypes_end0: | |
| } | |
| .section .debug_loc { } | |