| // | |
| // Generated by LLVM NVPTX Back-End | |
| // | |
| .version 8.2 | |
| .target sm_89 | |
| .address_size 64 | |
| // .globl triton__0d1d2d3d4d5d6de7de | |
| .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_1[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_1[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_1[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, 54, 32, 60, 32, 53, 48, 50, 53, 55}; | |
| .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[38] = {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, 51, 32, 60, 32, 53, 48, 50, 53, 55}; | |
| .global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0}; | |
| .visible .entry triton__0d1d2d3d4d5d6de7de( | |
| .param .u64 triton__0d1d2d3d4d5d6de7de_param_0, | |
| .param .u64 triton__0d1d2d3d4d5d6de7de_param_1, | |
| .param .u64 triton__0d1d2d3d4d5d6de7de_param_2, | |
| .param .u64 triton__0d1d2d3d4d5d6de7de_param_3, | |
| .param .u64 triton__0d1d2d3d4d5d6de7de_param_4, | |
| .param .u64 triton__0d1d2d3d4d5d6de7de_param_5, | |
| .param .u32 triton__0d1d2d3d4d5d6de7de_param_6, | |
| .param .u32 triton__0d1d2d3d4d5d6de7de_param_7 | |
| ) | |
| .maxntid 256, 1, 1 | |
| { | |
| .reg .pred %p<33>; | |
| .reg .b16 %rs<13>; | |
| .reg .b32 %r<93>; | |
| .reg .f32 %f<79>; | |
| .reg .b64 %rd<92>; | |
| .loc 1 18 0 | |
| $L__func_begin0: | |
| .loc 1 18 0 | |
| ld.param.u64 %rd37, [triton__0d1d2d3d4d5d6de7de_param_4]; | |
| ld.param.u64 %rd36, [triton__0d1d2d3d4d5d6de7de_param_3]; | |
| ld.param.u64 %rd35, [triton__0d1d2d3d4d5d6de7de_param_2]; | |
| ld.param.u64 %rd34, [triton__0d1d2d3d4d5d6de7de_param_1]; | |
| ld.param.u64 %rd43, [triton__0d1d2d3d4d5d6de7de_param_0]; | |
| $L__tmp0: | |
| .loc 1 22 44 | |
| mov.u32 %r1, %tid.x; | |
| bfe.u32 %r2, %r1, 2, 6; | |
| and.b32 %r16, %r1, 63; | |
| .loc 1 24 33 | |
| and.b32 %r3, %r1, 3; | |
| .loc 1 21 28 | |
| mov.u32 %r15, %ctaid.x; | |
| .loc 1 21 33 | |
| shl.b32 %r17, %r15, 6; | |
| .loc 1 22 23 | |
| or.b32 %r18, %r17, %r2; | |
| or.b32 %r19, %r17, %r16; | |
| .loc 1 26 30 | |
| mul.wide.s32 %rd44, %r18, 8; | |
| add.s64 %rd40, %rd43, %rd44; | |
| mul.wide.s32 %rd45, %r19, 8; | |
| add.s64 %rd42, %rd43, %rd45; | |
| mov.pred %p13, -1; | |
| .loc 1 26 35 | |
| mov.u64 %rd39, 0x0; | |
| @%p13 ld.global.L1::evict_last.b64 { %rd39 }, [ %rd40 + 0 ]; | |
| mov.u64 %rd41, 0x0; | |
| @%p13 ld.global.L1::evict_last.b64 { %rd41 }, [ %rd42 + 0 ]; | |
| .loc 1 27 18 | |
| bfe.s32 %r20, %r15, 25, 1; | |
| shr.u32 %r21, %r20, 23; | |
| add.s32 %r22, %r18, %r21; | |
| and.b32 %r23, %r22, 16776704; | |
| sub.s32 %r24, %r18, %r23; | |
| .loc 1 35 44 | |
| shl.b32 %r5, %r24, 8; | |
| .loc 1 37 22 | |
| add.s64 %rd46, %rd41, 50257; | |
| .loc 1 38 22 | |
| setp.lt.s64 %p3, %rd39, 0; | |
| setp.lt.s64 %p4, %rd41, 0; | |
| .loc 1 39 36 | |
| selp.b64 %rd47, %rd46, %rd41, %p4; | |
| .loc 1 40 40 | |
| setp.gt.u64 %p5, %rd47, 50256; | |
| .loc 1 41 44 | |
| shl.b64 %rd48, %rd39, 8; | |
| add.s64 %rd49, %rd48, 12865792; | |
| selp.b64 %rd2, %rd49, %rd48, %p3; | |
| mov.u16 %rs12, 0; | |
| mov.b32 %r76, 0; | |
| mov.b32 %r88, 883; | |
| mov.u64 %rd81, 1; | |
| .loc 1 40 55 | |
| @%p5 bra $L__BB0_3; | |
| bra.uni $L__BB0_1; | |
| $L__BB0_3: | |
| .loc 1 31 36 | |
| shl.b64 %rd55, %rd2, 2; | |
| mul.wide.u32 %rd88, %r3, 4; | |
| add.s64 %rd87, %rd55, %rd88; | |
| add.s64 %rd83, %rd34, %rd87; | |
| shl.b32 %r42, %r15, 14; | |
| shl.b32 %r43, %r2, 8; | |
| or.b32 %r44, %r42, %r43; | |
| or.b32 %r91, %r44, %r3; | |
| add.s32 %r45, %r5, %r3; | |
| mul.wide.s32 %rd86, %r45, 4; | |
| add.s64 %rd82, %rd35, %rd86; | |
| mov.f32 %f78, 0f00000000; | |
| mov.b32 %r89, -4; | |
| mov.f32 %f77, %f78; | |
| mov.f32 %f76, %f78; | |
| $L__BB0_4: | |
| .loc 1 35 50 | |
| mov.u32 %r46, 0x0; | |
| @%p13 ld.global.L1::evict_last.b32 { %r46 }, [ %rd82 + 0 ]; | |
| @!%p13 mov.u32 %r46, %r76; | |
| mov.b32 %f31, %r46; | |
| .loc 1 31 36 | |
| add.s32 %r89, %r89, 4; | |
| .loc 1 36 34 | |
| add.s32 %r54, %r89, %r91; | |
| mul.wide.s32 %rd59, %r54, 2; | |
| add.s64 %rd57, %rd36, %rd59; | |
| .loc 1 36 50 | |
| mov.u16 %rs4, 0x0; | |
| @%p13 ld.global.L1::evict_last.b16 { %rs4 }, [ %rd57 + 0 ]; | |
| @!%p13 mov.u16 %rs4, %rs12; | |
| .loc 1 36 101 | |
| cvt.f32.bf16 %r48, %rs4; | |
| mov.b32 %f32, %r48; | |
| .loc 1 40 55 | |
| mov.u64 %rd60, assertMessage_0; | |
| cvta.global.u64 %rd61, %rd60; | |
| mov.u64 %rd62, assertFile_0; | |
| cvta.global.u64 %rd63, %rd62; | |
| mov.u64 %rd64, assertFunc_0; | |
| cvta.global.u64 %rd65, %rd64; | |
| { // callseq 10, 0 | |
| .reg .b32 temp_param_reg; | |
| .param .b64 param0; | |
| st.param.b64 [param0+0], %rd61; | |
| .param .b64 param1; | |
| st.param.b64 [param1+0], %rd63; | |
| .param .b32 param2; | |
| st.param.b32 [param2+0], %r88; | |
| .param .b64 param3; | |
| st.param.b64 [param3+0], %rd65; | |
| .param .b64 param4; | |
| st.param.b64 [param4+0], %rd81; | |
| call.uni | |
| __assertfail, | |
| ( | |
| param0, | |
| param1, | |
| param2, | |
| param3, | |
| param4 | |
| ); | |
| } // callseq 10 | |
| .loc 1 41 52 | |
| mov.u32 %r49, 0x0; | |
| @%p13 ld.global.L1::evict_last.b32 { %r49 }, [ %rd83 + 0 ]; | |
| @!%p13 mov.u32 %r49, %r76; | |
| mov.b32 %f33, %r49; | |
| .loc 1 42 22 | |
| add.f32 %f34, %f31, %f33; | |
| .loc 1 44 22 | |
| add.f32 %f35, %f32, %f34; | |
| $L__tmp1: | |
| .loc 2 96 20 | |
| sub.f32 %f36, %f35, %f76; | |
| .loc 2 97 26 | |
| add.f32 %f78, %f78, 0f3F800000; | |
| .loc 2 98 30 | |
| mov.b32 %r52, %f36; | |
| mov.b32 %r53, %f78; | |
| div.full.f32 %r51, %r52, %r53; | |
| mov.b32 %f37, %r51; | |
| .loc 2 98 22 | |
| add.f32 %f76, %f76, %f37; | |
| .loc 2 101 30 | |
| sub.f32 %f38, %f35, %f76; | |
| $L__tmp2: | |
| .loc 1 50 50 | |
| fma.rn.f32 %f77, %f36, %f38, %f77; | |
| .loc 1 31 36 | |
| add.s64 %rd83, %rd83, 16; | |
| add.s64 %rd82, %rd82, 16; | |
| setp.lt.u32 %p19, %r89, 252; | |
| @%p19 bra $L__BB0_4; | |
| bra.uni $L__BB0_5; | |
| $L__BB0_1: | |
| .loc 1 0 36 | |
| mov.b32 %r90, -4; | |
| .loc 1 31 36 | |
| shl.b64 %rd50, %rd2, 2; | |
| mul.wide.u32 %rd88, %r3, 4; | |
| add.s64 %rd87, %rd50, %rd88; | |
| add.s64 %rd85, %rd34, %rd87; | |
| shl.b32 %r27, %r15, 14; | |
| shl.b32 %r28, %r2, 8; | |
| or.b32 %r29, %r27, %r28; | |
| or.b32 %r91, %r29, %r3; | |
| add.s32 %r30, %r5, %r3; | |
| mul.wide.s32 %rd86, %r30, 4; | |
| add.s64 %rd84, %rd35, %rd86; | |
| mov.f32 %f78, 0f00000000; | |
| mov.f32 %f77, %f78; | |
| mov.f32 %f76, %f78; | |
| $L__BB0_2: | |
| .loc 1 35 50 | |
| mov.u32 %r31, 0x0; | |
| @%p13 ld.global.L1::evict_last.b32 { %r31 }, [ %rd84 + 0 ]; | |
| @!%p13 mov.u32 %r31, %r76; | |
| mov.b32 %f22, %r31; | |
| .loc 1 31 36 | |
| add.s32 %r90, %r90, 4; | |
| .loc 1 36 34 | |
| add.s32 %r39, %r90, %r91; | |
| mul.wide.s32 %rd54, %r39, 2; | |
| add.s64 %rd52, %rd36, %rd54; | |
| .loc 1 36 50 | |
| mov.u16 %rs1, 0x0; | |
| @%p13 ld.global.L1::evict_last.b16 { %rs1 }, [ %rd52 + 0 ]; | |
| @!%p13 mov.u16 %rs1, %rs12; | |
| .loc 1 36 101 | |
| cvt.f32.bf16 %r33, %rs1; | |
| mov.b32 %f23, %r33; | |
| .loc 1 41 52 | |
| mov.u32 %r34, 0x0; | |
| @%p13 ld.global.L1::evict_last.b32 { %r34 }, [ %rd85 + 0 ]; | |
| @!%p13 mov.u32 %r34, %r76; | |
| mov.b32 %f24, %r34; | |
| .loc 1 42 22 | |
| add.f32 %f25, %f22, %f24; | |
| .loc 1 44 22 | |
| add.f32 %f26, %f23, %f25; | |
| $L__tmp3: | |
| .loc 2 96 20 | |
| sub.f32 %f27, %f26, %f76; | |
| .loc 2 97 26 | |
| add.f32 %f78, %f78, 0f3F800000; | |
| .loc 2 98 30 | |
| mov.b32 %r37, %f27; | |
| mov.b32 %r38, %f78; | |
| div.full.f32 %r36, %r37, %r38; | |
| mov.b32 %f28, %r36; | |
| .loc 2 98 22 | |
| add.f32 %f76, %f76, %f28; | |
| .loc 2 101 30 | |
| sub.f32 %f29, %f26, %f76; | |
| $L__tmp4: | |
| .loc 1 50 50 | |
| fma.rn.f32 %f77, %f27, %f29, %f77; | |
| .loc 1 31 36 | |
| add.s64 %rd85, %rd85, 16; | |
| add.s64 %rd84, %rd84, 16; | |
| setp.lt.u32 %p12, %r90, 252; | |
| @%p12 bra $L__BB0_2; | |
| $L__BB0_5: | |
| .loc 1 0 36 | |
| ld.param.u64 %rd38, [triton__0d1d2d3d4d5d6de7de_param_5]; | |
| $L__tmp5: | |
| .loc 2 120 46 | |
| mov.b32 %r66, %f76; | |
| shfl.sync.bfly.b32 %r67, %r66, 2, 31, -1; | |
| mov.b32 %f39, %r67; | |
| mov.b32 %r68, %f77; | |
| shfl.sync.bfly.b32 %r69, %r68, 2, 31, -1; | |
| mov.b32 %f40, %r69; | |
| mov.b32 %r70, %f78; | |
| shfl.sync.bfly.b32 %r57, %r70, 2, 31, -1; | |
| mov.b32 %f41, %r57; | |
| $L__tmp6: | |
| .loc 2 108 21 | |
| sub.f32 %f42, %f39, %f76; | |
| .loc 2 109 28 | |
| add.f32 %f43, %f78, %f41; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p20, %f43, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r58, %f43; | |
| div.full.f32 %r56, %r57, %r58; | |
| mov.b32 %f44, %r56; | |
| .loc 2 110 49 | |
| selp.f32 %f45, 0f00000000, %f44, %p20; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f46, %f42, %f45, %f76; | |
| .loc 2 113 15 | |
| add.f32 %f47, %f77, %f40; | |
| .loc 2 113 30 | |
| mul.f32 %f48, %f42, %f42; | |
| .loc 2 113 38 | |
| mul.f32 %f49, %f78, %f48; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f50, %f49, %f45, %f47; | |
| $L__tmp7: | |
| .loc 2 120 46 | |
| mov.b32 %r71, %f46; | |
| shfl.sync.bfly.b32 %r72, %r71, 1, 31, -1; | |
| mov.b32 %f51, %r72; | |
| mov.b32 %r73, %f50; | |
| shfl.sync.bfly.b32 %r74, %r73, 1, 31, -1; | |
| mov.b32 %f52, %r74; | |
| shfl.sync.bfly.b32 %r60, %r58, 1, 31, -1; | |
| mov.b32 %f53, %r60; | |
| $L__tmp8: | |
| .loc 2 108 21 | |
| sub.f32 %f54, %f51, %f46; | |
| .loc 2 109 28 | |
| add.f32 %f55, %f43, %f53; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p21, %f55, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r61, %f55; | |
| div.full.f32 %r59, %r60, %r61; | |
| mov.b32 %f56, %r59; | |
| .loc 2 110 49 | |
| selp.f32 %f57, 0f00000000, %f56, %p21; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f16, %f54, %f57, %f46; | |
| .loc 2 113 15 | |
| add.f32 %f58, %f50, %f52; | |
| .loc 2 113 30 | |
| mul.f32 %f59, %f54, %f54; | |
| .loc 2 113 38 | |
| mul.f32 %f60, %f43, %f59; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f61, %f57, %f60, %f58; | |
| $L__tmp9: | |
| .loc 1 75 24 | |
| mov.b32 %r63, %f61; | |
| mov.b32 %r64, 1132462080; | |
| div.full.f32 %r62, %r63, %r64; | |
| mov.b32 %f62, %r62; | |
| .loc 1 77 24 | |
| add.f32 %f17, %f62, 0f3727C5AC; | |
| .loc 1 58 36 | |
| add.s64 %rd91, %rd34, %rd87; | |
| add.s64 %rd90, %rd37, %rd88; | |
| add.s64 %rd89, %rd35, %rd86; | |
| mov.b32 %r92, -4; | |
| setp.lt.u64 %p28, %rd47, 50257; | |
| rsqrt.approx.ftz.f32 %f67, %f17; | |
| bra.uni $L__BB0_6; | |
| $L__BB0_8: | |
| .loc 1 0 0 | |
| mov.b32 %f18, %r75; | |
| cvt.s64.s32 %rd30, %r81; | |
| cvt.f32.bf16 %r77, %rs7; | |
| mov.b32 %f19, %r77; | |
| mov.b32 %f20, %r78; | |
| .loc 1 69 54 | |
| mov.u32 %r83, 0x0; | |
| @%p13 ld.global.L1::evict_first.b32 { %r83 }, [ %rd91 + 0 ]; | |
| @!%p13 mov.u32 %r83, %r76; | |
| mov.b32 %f63, %r83; | |
| .loc 1 70 24 | |
| add.f32 %f64, %f18, %f63; | |
| .loc 1 72 24 | |
| add.f32 %f65, %f19, %f64; | |
| .loc 1 73 24 | |
| sub.f32 %f66, %f65, %f16; | |
| .loc 1 79 24 | |
| mul.f32 %f68, %f66, %f67; | |
| .loc 1 80 24 | |
| mul.f32 %f69, %f68, %f20; | |
| .loc 1 82 29 | |
| shl.b64 %rd80, %rd30, 1; | |
| add.s64 %rd79, %rd38, %rd80; | |
| .loc 1 82 52 | |
| mov.b32 %r85, %f69; | |
| cvt.rn.bf16.f32 %rs10, %r85; | |
| @%p13 st.global.b16 [ %rd79 + 0 ], { %rs10 }; | |
| .loc 1 58 36 | |
| add.s32 %r92, %r92, 4; | |
| add.s64 %rd91, %rd91, 16; | |
| add.s64 %rd90, %rd90, 16; | |
| add.s64 %rd89, %rd89, 16; | |
| setp.lt.u32 %p32, %r92, 252; | |
| @%p32 bra $L__BB0_6; | |
| bra.uni $L__BB0_9; | |
| $L__BB0_6: | |
| .loc 1 62 51 | |
| mov.u32 %r75, 0x0; | |
| @%p13 ld.global.L1::evict_last.b32 { %r75 }, [ %rd89 + 0 ]; | |
| @!%p13 mov.u32 %r75, %r76; | |
| .loc 1 63 35 | |
| add.s32 %r80, %r91, %r92; | |
| add.s32 %r81, %r80, 4; | |
| mul.wide.s32 %rd70, %r81, 2; | |
| add.s64 %rd68, %rd36, %rd70; | |
| .loc 1 63 51 | |
| mov.u16 %rs7, 0x0; | |
| @%p13 ld.global.L1::evict_first.b16 { %rs7 }, [ %rd68 + 0 ]; | |
| @!%p13 mov.u16 %rs7, %rs12; | |
| .loc 1 64 40 | |
| mov.u32 %r78, 0x0; | |
| @%p13 ld.global.L1::evict_last.b32 { %r78 }, [ %rd90 + 0 ]; | |
| @!%p13 mov.u32 %r78, %r76; | |
| .loc 1 68 57 | |
| @%p28 bra $L__BB0_8; | |
| mov.u64 %rd71, assertMessage_1; | |
| cvta.global.u64 %rd72, %rd71; | |
| mov.u64 %rd73, assertFile_1; | |
| cvta.global.u64 %rd74, %rd73; | |
| mov.u64 %rd75, assertFunc_1; | |
| cvta.global.u64 %rd76, %rd75; | |
| { // callseq 11, 0 | |
| .reg .b32 temp_param_reg; | |
| .param .b64 param0; | |
| st.param.b64 [param0+0], %rd72; | |
| .param .b64 param1; | |
| st.param.b64 [param1+0], %rd74; | |
| .param .b32 param2; | |
| st.param.b32 [param2+0], %r88; | |
| .param .b64 param3; | |
| st.param.b64 [param3+0], %rd76; | |
| .param .b64 param4; | |
| st.param.b64 [param4+0], %rd81; | |
| call.uni | |
| __assertfail, | |
| ( | |
| param0, | |
| param1, | |
| param2, | |
| param3, | |
| param4 | |
| ); | |
| } // callseq 11 | |
| bra.uni $L__BB0_8; | |
| $L__BB0_9: | |
| .loc 1 58 4 | |
| ret; | |
| $L__tmp10: | |
| $L__func_end0: | |
| } | |
| // .globl __nv_rsqrtf | |
| .visible .func (.param .b32 func_retval0) __nv_rsqrtf( | |
| .param .b32 __nv_rsqrtf_param_0 | |
| ) | |
| { | |
| .reg .f32 %f<3>; | |
| $L__func_begin1: | |
| ld.param.f32 %f1, [__nv_rsqrtf_param_0]; | |
| rsqrt.approx.ftz.f32 %f2, %f1; | |
| st.param.f32 [func_retval0+0], %f2; | |
| ret; | |
| $L__func_end1: | |
| } | |
| .file 1 "/tmp/torchinductor_root/ci/ccig6fki6p4lxrdmgg6eudahiexcvueeol2p4qp532pvve2y463y.py" | |
| .file 2 "/usr/local/lib/python3.10/dist-packages/torch/_inductor/triton_helpers.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 302 | |
| .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 99 | |
| .b8 105 | |
| .b8 103 | |
| .b8 54 | |
| .b8 102 | |
| .b8 107 | |
| .b8 105 | |
| .b8 54 | |
| .b8 112 | |
| .b8 52 | |
| .b8 108 | |
| .b8 120 | |
| .b8 114 | |
| .b8 100 | |
| .b8 109 | |
| .b8 103 | |
| .b8 103 | |
| .b8 54 | |
| .b8 101 | |
| .b8 117 | |
| .b8 100 | |
| .b8 97 | |
| .b8 104 | |
| .b8 105 | |
| .b8 101 | |
| .b8 120 | |
| .b8 99 | |
| .b8 118 | |
| .b8 117 | |
| .b8 101 | |
| .b8 101 | |
| .b8 111 | |
| .b8 108 | |
| .b8 50 | |
| .b8 112 | |
| .b8 52 | |
| .b8 113 | |
| .b8 112 | |
| .b8 53 | |
| .b8 51 | |
| .b8 50 | |
| .b8 112 | |
| .b8 118 | |
| .b8 118 | |
| .b8 101 | |
| .b8 50 | |
| .b8 121 | |
| .b8 52 | |
| .b8 54 | |
| .b8 51 | |
| .b8 121 | |
| .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 99 | |
| .b8 105 | |
| .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 100 | |
| .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 100 | |
| .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__tmp4 | |
| .b8 2 | |
| .b8 47 | |
| .b8 41 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp5 | |
| .b64 $L__tmp8 | |
| .b8 2 | |
| .b8 53 | |
| .b8 44 | |
| .b8 5 | |
| .b32 125 | |
| .b64 $L__tmp6 | |
| .b64 $L__tmp9 | |
| .b8 2 | |
| .b8 53 | |
| .b8 44 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp6 | |
| .b64 $L__tmp9 | |
| .b8 2 | |
| .b8 120 | |
| .b8 46 | |
| .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 306 | |
| .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 100 | |
| .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 306 | |
| .b32 0 | |
| $L__pubTypes_end0: | |
| } | |
| .section .debug_loc { } | |