// // Generated by LLVM NVPTX Back-End // .version 8.8 .target sm_89 .address_size 64 // .globl _attn_fwd // -- Begin function _attn_fwd .extern .shared .align 16 .b8 global_smem[]; .global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90}; // @_attn_fwd .visible .entry _attn_fwd( .param .u64 .ptr .global .align 1 _attn_fwd_param_0, .param .u64 .ptr .global .align 1 _attn_fwd_param_1, .param .u64 .ptr .global .align 1 _attn_fwd_param_2, .param .u64 .ptr .global .align 1 _attn_fwd_param_3, .param .u64 .ptr .global .align 1 _attn_fwd_param_4, .param .u64 .ptr .global .align 1 _attn_fwd_param_5, .param .u32 _attn_fwd_param_6, .param .u32 _attn_fwd_param_7, .param .u32 _attn_fwd_param_8, .param .u32 _attn_fwd_param_9, .param .u32 _attn_fwd_param_10, .param .u32 _attn_fwd_param_11, .param .u32 _attn_fwd_param_12, .param .u32 _attn_fwd_param_13, .param .u32 _attn_fwd_param_14, .param .u32 _attn_fwd_param_15, .param .u32 _attn_fwd_param_16, .param .u32 _attn_fwd_param_17, .param .u32 _attn_fwd_param_18, .param .u32 _attn_fwd_param_19, .param .u64 .ptr .global .align 1 _attn_fwd_param_20, .param .u64 .ptr .global .align 1 _attn_fwd_param_21 ) .reqntid 128 { .reg .pred %p<39>; .reg .b16 %rs<65>; .reg .b32 %r<1246>; .reg .b64 %rd<128>; .loc 1 40 0 // attn_qk_int8_per_block.py:40:0 $L__func_begin0: .loc 1 40 0 // attn_qk_int8_per_block.py:40:0 // %bb.0: ld.param.b32 %r38, [_attn_fwd_param_19]; ld.param.b32 %r37, [_attn_fwd_param_18]; ld.param.b32 %r36, [_attn_fwd_param_17]; ld.param.b32 %r35, [_attn_fwd_param_16]; ld.param.b32 %r34, [_attn_fwd_param_15]; ld.param.b64 %rd26, [_attn_fwd_param_5]; ld.param.b64 %rd25, [_attn_fwd_param_4]; ld.param.b64 %rd24, [_attn_fwd_param_2]; ld.param.b64 %rd23, [_attn_fwd_param_1]; ld.param.b64 %rd46, [_attn_fwd_param_0]; $L__tmp0: .loc 1 51 15 // attn_qk_int8_per_block.py:51:15 mov.u32 %r88, %ctaid.x; .loc 1 53 13 // attn_qk_int8_per_block.py:53:13 mov.u32 %r89, %ctaid.z; cvt.u64.u32 %rd1, %r89; ld.param.b64 %rd47, [_attn_fwd_param_3]; .loc 1 54 13 // attn_qk_int8_per_block.py:54:13 mov.u32 %r90, %ctaid.y; cvt.u64.u32 %rd2, %r90; ld.param.s32 %rd48, [_attn_fwd_param_6]; .loc 1 56 23 // attn_qk_int8_per_block.py:56:23 mad.wide.u32 %rd49, %r89, 32, %rd2; ld.param.s32 %rd50, [_attn_fwd_param_7]; ld.param.b32 %r91, [_attn_fwd_param_8]; ld.param.s32 %rd51, [_attn_fwd_param_9]; $L__tmp1: .loc 1 56 44 // attn_qk_int8_per_block.py:56:44 .loc 2 43 13, function_name $L__info_string0, inlined_at 1 56 44 // standard.py:43:13 add.s32 %r92, %r37, 127; add.s32 %r93, %r38, 63; ld.param.s32 %rd52, [_attn_fwd_param_10]; .loc 2 43 12, function_name $L__info_string0, inlined_at 1 56 44 // standard.py:43:12 shr.s32 %r94, %r93, 31; shr.u32 %r95, %r94, 26; add.s32 %r96, %r93, %r95; shr.s32 %r97, %r96, 6; shr.s32 %r98, %r92, 31; shr.u32 %r99, %r98, 25; add.s32 %r100, %r92, %r99; shr.s32 %r101, %r100, 7; ld.param.b32 %r102, [_attn_fwd_param_11]; ld.param.s32 %rd53, [_attn_fwd_param_12]; $L__tmp2: .loc 1 56 22 // attn_qk_int8_per_block.py:56:22 cvt.s64.s32 %rd54, %r101; ld.param.s32 %rd55, [_attn_fwd_param_13]; mul.lo.s64 %rd56, %rd49, %rd54; ld.param.b32 %r103, [_attn_fwd_param_14]; .loc 1 57 22 // attn_qk_int8_per_block.py:57:22 cvt.s64.s32 %rd57, %r97; mul.lo.s64 %rd3, %rd49, %rd57; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 shl.b32 %r1, %r88, 7; .loc 1 59 34 // attn_qk_int8_per_block.py:59:34 mov.u32 %r2, %tid.x; bfe.u32 %r3, %r2, 2, 5; or.b32 %r4, %r3, 32; shr.u32 %r5, %r2, 3; bfe.u32 %r6, %r2, 3, 4; or.b32 %r7, %r6, 16; or.b32 %r8, %r6, 32; or.b32 %r9, %r6, 48; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r104, %r3, %r1; or.b32 %r105, %r4, %r1; or.b32 %r106, %r104, 64; or.b32 %r107, %r104, 96; .loc 1 62 14 // attn_qk_int8_per_block.py:62:14 mad.lo.s64 %rd58, %rd48, %rd1, %rd46; mad.lo.s64 %rd59, %rd50, %rd2, %rd58; .loc 1 62 60 // attn_qk_int8_per_block.py:62:60 mul.lo.s32 %r108, %r91, %r104; mul.lo.s32 %r109, %r91, %r105; mul.lo.s32 %r110, %r91, %r106; mul.lo.s32 %r111, %r91, %r107; .loc 1 62 14 // attn_qk_int8_per_block.py:62:14 cvt.s64.s32 %rd60, %r108; add.s64 %rd61, %rd59, %rd60; cvt.s64.s32 %rd62, %r109; add.s64 %rd63, %rd59, %rd62; cvt.s64.s32 %rd64, %r110; add.s64 %rd65, %rd59, %rd64; cvt.s64.s32 %rd66, %r111; add.s64 %rd67, %rd59, %rd66; .loc 1 62 90 // attn_qk_int8_per_block.py:62:90 shl.b32 %r10, %r2, 4; and.b32 %r112, %r10, 48; and.b32 %r11, %r2, 7; shl.b32 %r113, %r11, 3; .loc 1 62 14 // attn_qk_int8_per_block.py:62:14 cvt.u64.u32 %rd4, %r112; add.s64 %rd27, %rd61, %rd4; add.s64 %rd28, %rd63, %rd4; add.s64 %rd29, %rd65, %rd4; add.s64 %rd30, %rd67, %rd4; .loc 1 63 19 // attn_qk_int8_per_block.py:63:19 shl.b64 %rd68, %rd56, 2; add.s64 %rd69, %rd47, %rd68; mad.wide.u32 %rd31, %r88, 4, %rd69; .loc 1 64 19 // attn_qk_int8_per_block.py:64:19 mul.lo.s64 %rd5, %rd51, %rd1; .loc 1 64 39 // attn_qk_int8_per_block.py:64:39 mul.lo.s64 %rd6, %rd52, %rd2; .loc 1 64 14 // attn_qk_int8_per_block.py:64:14 add.s64 %rd70, %rd23, %rd5; add.s64 %rd71, %rd70, %rd6; .loc 1 64 79 // attn_qk_int8_per_block.py:64:79 mul.lo.s32 %r114, %r102, %r3; shl.b32 %r115, %r102, 5; add.s32 %r116, %r114, %r115; .loc 1 64 14 // attn_qk_int8_per_block.py:64:14 cvt.s64.s32 %rd7, %r114; add.s64 %rd72, %rd71, %rd7; cvt.s64.s32 %rd8, %r116; add.s64 %rd73, %rd71, %rd8; add.s64 %rd32, %rd72, %rd4; add.s64 %rd33, %rd73, %rd4; .loc 1 65 19 // attn_qk_int8_per_block.py:65:19 shl.b64 %rd74, %rd3, 2; add.s64 %rd34, %rd25, %rd74; .loc 1 66 19 // attn_qk_int8_per_block.py:66:19 mul.lo.s64 %rd9, %rd53, %rd1; .loc 1 66 39 // attn_qk_int8_per_block.py:66:39 mul.lo.s64 %rd10, %rd55, %rd2; .loc 1 66 14 // attn_qk_int8_per_block.py:66:14 shl.b64 %rd75, %rd9, 1; add.s64 %rd76, %rd24, %rd75; shl.b64 %rd77, %rd10, 1; add.s64 %rd78, %rd76, %rd77; .loc 1 66 79 // attn_qk_int8_per_block.py:66:79 mul.lo.s32 %r117, %r103, %r6; shl.b32 %r118, %r103, 4; add.s32 %r119, %r117, %r118; shl.b32 %r120, %r103, 5; add.s32 %r121, %r117, %r120; mad.lo.s32 %r122, %r103, 48, %r117; .loc 1 66 14 // attn_qk_int8_per_block.py:66:14 mad.wide.s32 %rd79, %r117, 2, %rd78; mad.wide.s32 %rd80, %r119, 2, %rd78; mad.wide.s32 %rd81, %r121, 2, %rd78; mad.wide.s32 %rd82, %r122, 2, %rd78; cvt.u64.u32 %rd15, %r113; mul.wide.u32 %rd83, %r113, 2; add.s64 %rd35, %rd79, %rd83; add.s64 %rd36, %rd80, %rd83; add.s64 %rd37, %rd81, %rd83; add.s64 %rd38, %rd82, %rd83; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p1, %r104, %r37; setp.lt.s32 %p2, %r105, %r37; setp.lt.s32 %p3, %r106, %r37; setp.lt.s32 %p4, %r107, %r37; .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r39, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r40, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r41, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r42, 0x0; @%p1 ld.global.v4.b32 { %r39, %r40, %r41, %r42 }, [ %rd27 + 0 ]; // end inline asm // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r43, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r44, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r45, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r46, 0x0; @%p2 ld.global.v4.b32 { %r43, %r44, %r45, %r46 }, [ %rd28 + 0 ]; // end inline asm // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r47, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r48, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r49, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r50, 0x0; @%p3 ld.global.v4.b32 { %r47, %r48, %r49, %r50 }, [ %rd29 + 0 ]; // end inline asm // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r51, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r52, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r53, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r54, 0x0; @%p4 ld.global.v4.b32 { %r51, %r52, %r53, %r54 }, [ %rd30 + 0 ]; // end inline asm and.b32 %r123, %r10, 2032; shl.b32 %r124, %r2, 1; and.b32 %r12, %r124, 48; xor.b32 %r125, %r123, %r12; mov.b32 %r126, global_smem; add.s32 %r13, %r126, %r125; add.s32 %r55, %r13, 16384; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r55 + 0 ], { %r39, %r40, %r41, %r42 }; // end inline asm add.s32 %r56, %r13, 18432; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r56 + 0 ], { %r43, %r44, %r45, %r46 }; // end inline asm add.s32 %r57, %r13, 20480; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r57 + 0 ], { %r47, %r48, %r49, %r50 }; // end inline asm add.s32 %r58, %r13, 22528; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r58 + 0 ], { %r51, %r52, %r53, %r54 }; // end inline asm .loc 1 74 15 // attn_qk_int8_per_block.py:74:15 // begin inline asm .loc 1 74 15 // attn_qk_int8_per_block.py:74:15 mov.u32 %r59, 0x0; .loc 1 75 15 // attn_qk_int8_per_block.py:75:15 ld.global.b32 { %r59 }, [ %rd31 + 0 ]; // end inline asm $L__tmp3: .loc 1 75 16 // attn_qk_int8_per_block.py:75:16 .loc 1 34 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:34:19 shl.b32 %r127, %r102, 6; .loc 1 36 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:36:19 shl.b32 %r128, %r103, 6; .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 setp.gt.s32 %p6, %r38, 0; .loc 1 15 18, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:15:18 setp.lt.s32 %p7, %r3, %r38; setp.lt.s32 %p8, %r4, %r38; .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 add.s32 %r60, %r13, 24576; selp.b32 %r61, 16, 0, %p7; // begin inline asm .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.cg.shared.global [ %r60 + 0 ], [ %rd32 + 0 ], 0x10, %r61; // end inline asm add.s32 %r62, %r13, 26624; selp.b32 %r129, 16, 0, %p8; selp.b32 %r63, %r129, 0, %p6; // begin inline asm .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.cg.shared.global [ %r62 + 0 ], [ %rd33 + 0 ], 0x10, %r63; // end inline asm cp.async.commit_group; .loc 1 17 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:17:19 and.b32 %r14, %r2, 127; setp.eq.b32 %p5, %r14, 0; selp.b32 %r65, 4, 0, %p6; add.s32 %r64, %r126, 32768; // begin inline asm @%p5 cp.async.ca.shared.global [ %r64 + 0 ], [ %rd34 + 0 ], 0x4, %r65; // end inline asm cp.async.commit_group; .loc 1 29 36, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:36 setp.lt.s32 %p9, %r6, %r38; setp.lt.s32 %p10, %r7, %r38; setp.lt.s32 %p11, %r8, %r38; setp.lt.s32 %p12, %r9, %r38; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 and.b32 %r130, %r124, 112; xor.b32 %r15, %r123, %r130; add.s32 %r66, %r126, %r15; selp.b32 %r67, 16, 0, %p9; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r66 + 0 ], [ %rd35 + 0 ], 0x10, %r67; // end inline asm add.s32 %r68, %r66, 2048; selp.b32 %r131, 16, 0, %p10; selp.b32 %r69, %r131, 0, %p6; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r68 + 0 ], [ %rd36 + 0 ], 0x10, %r69; // end inline asm add.s32 %r70, %r66, 4096; selp.b32 %r132, 16, 0, %p11; selp.b32 %r71, %r132, 0, %p6; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r70 + 0 ], [ %rd37 + 0 ], 0x10, %r71; // end inline asm add.s32 %r72, %r66, 6144; selp.b32 %r133, 16, 0, %p12; selp.b32 %r73, %r133, 0, %p6; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r72 + 0 ], [ %rd38 + 0 ], 0x10, %r73; // end inline asm cp.async.commit_group; .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 setp.gt.s32 %p13, %r38, 64; .loc 1 34 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:34:9 cvt.s64.s32 %rd16, %r127; add.s64 %rd39, %rd32, %rd16; add.s64 %rd40, %rd33, %rd16; .loc 1 35 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:35:9 add.s64 %rd41, %rd34, 4; .loc 1 36 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:36:9 mul.wide.s32 %rd84, %r128, 2; add.s64 %rd42, %rd35, %rd84; add.s64 %rd43, %rd36, %rd84; add.s64 %rd44, %rd37, %rd84; add.s64 %rd45, %rd38, %rd84; .loc 1 15 37, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:15:37 add.s32 %r134, %r38, -64; .loc 1 15 18, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:15:18 setp.lt.s32 %p14, %r3, %r134; setp.lt.s32 %p15, %r4, %r134; .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 bar.sync 0; add.s32 %r74, %r13, 28672; selp.b32 %r135, 16, 0, %p14; selp.b32 %r75, %r135, 0, %p13; // begin inline asm .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.cg.shared.global [ %r74 + 0 ], [ %rd39 + 0 ], 0x10, %r75; // end inline asm add.s32 %r76, %r13, 30720; selp.b32 %r136, 16, 0, %p15; selp.b32 %r77, %r136, 0, %p13; // begin inline asm .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.cg.shared.global [ %r76 + 0 ], [ %rd40 + 0 ], 0x10, %r77; // end inline asm cp.async.commit_group; .loc 1 17 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:17:19 selp.b32 %r79, 4, 0, %p13; add.s32 %r78, %r126, 32772; // begin inline asm @%p5 cp.async.ca.shared.global [ %r78 + 0 ], [ %rd41 + 0 ], 0x4, %r79; // end inline asm cp.async.commit_group; .loc 1 29 36, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:36 setp.lt.s32 %p16, %r6, %r134; setp.lt.s32 %p17, %r7, %r134; setp.lt.s32 %p18, %r8, %r134; setp.lt.s32 %p19, %r9, %r134; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 add.s32 %r80, %r66, 8192; selp.b32 %r137, 16, 0, %p16; selp.b32 %r81, %r137, 0, %p13; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r80 + 0 ], [ %rd42 + 0 ], 0x10, %r81; // end inline asm add.s32 %r82, %r66, 10240; selp.b32 %r138, 16, 0, %p17; selp.b32 %r83, %r138, 0, %p13; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r82 + 0 ], [ %rd43 + 0 ], 0x10, %r83; // end inline asm add.s32 %r84, %r66, 12288; selp.b32 %r139, 16, 0, %p18; selp.b32 %r85, %r139, 0, %p13; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r84 + 0 ], [ %rd44 + 0 ], 0x10, %r85; // end inline asm add.s32 %r86, %r66, 14336; selp.b32 %r140, 16, 0, %p19; selp.b32 %r87, %r140, 0, %p13; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r86 + 0 ], [ %rd45 + 0 ], 0x10, %r87; // end inline asm cp.async.commit_group; .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 @!%p6 bra $L__BB0_1; // %bb.2: // %.lr.ph .loc 1 0 5 // attn_qk_int8_per_block.py:0:5 cvt.s64.s32 %rd11, %r117; cvt.s64.s32 %rd12, %r119; cvt.s64.s32 %rd13, %r121; cvt.s64.s32 %rd14, %r122; cvt.s64.s32 %rd17, %r128; add.s32 %r16, %r38, -128; shl.b32 %r141, %r2, 6; and.b32 %r142, %r141, 960; shl.b32 %r143, %r2, 5; and.b32 %r144, %r143, 3072; shl.b32 %r145, %r2, 3; and.b32 %r146, %r145, 48; and.b32 %r1177, %r2, 16; or.b32 %r147, %r142, %r146; xor.b32 %r148, %r147, %r1177; or.b32 %r149, %r148, %r144; add.s32 %r150, %r126, %r149; add.s32 %r17, %r150, 16384; add.s32 %r18, %r150, 20480; xor.b32 %r151, %r149, 32; add.s32 %r152, %r126, %r151; add.s32 %r19, %r152, 16384; add.s32 %r20, %r152, 20480; shl.b32 %r153, %r11, 6; or.b32 %r154, %r153, %r146; xor.b32 %r155, %r154, %r12; .loc 1 13 5 // attn_qk_int8_per_block.py:13:5 @[ attn_qk_int8_per_block.py:75:16 ] add.s32 %r21, %r126, %r155; shl.b32 %r156, %r2, 7; and.b32 %r157, %r156, 3968; shl.b32 %r158, %r11, 4; or.b32 %r22, %r158, %r157; xor.b32 %r23, %r22, 16; xor.b32 %r24, %r22, 32; xor.b32 %r25, %r22, 48; xor.b32 %r26, %r22, 64; xor.b32 %r27, %r22, 80; xor.b32 %r28, %r22, 96; xor.b32 %r29, %r22, 112; shl.b64 %rd18, %rd14, 1; add.s64 %rd85, %rd77, %rd75; mad.wide.u32 %rd86, %r11, 16, %rd85; shl.b64 %rd87, %rd17, 2; add.s64 %rd88, %rd86, %rd87; add.s64 %rd127, %rd24, %rd88; shl.b64 %rd19, %rd17, 1; shl.b64 %rd20, %rd13, 1; shl.b64 %rd21, %rd12, 1; shl.b64 %rd22, %rd11, 1; add.s64 %rd89, %rd6, %rd5; add.s64 %rd90, %rd89, %rd4; shl.b64 %rd91, %rd16, 1; add.s64 %rd92, %rd90, %rd91; add.s64 %rd126, %rd23, %rd92; add.s64 %rd93, %rd74, %rd25; add.s64 %rd125, %rd93, 8; mov.b32 %r1173, 0fFF800000; mov.b32 %r1242, 0f3F800000; mov.b32 %r1178, 0f00000000; mov.b32 %r159, 0; mov.b32 %r1171, 1; mov.b32 %r1170, -1; mov.b32 %r1169, %r16; mov.b32 %r1172, %r159; mov.b32 %r1179, %r1178; mov.b32 %r1180, %r1178; mov.b32 %r1181, %r1178; mov.b32 %r1182, %r1178; mov.b32 %r1183, %r1178; mov.b32 %r1184, %r1178; mov.b32 %r1185, %r1178; mov.b32 %r1186, %r1178; mov.b32 %r1187, %r1178; mov.b32 %r1188, %r1178; mov.b32 %r1189, %r1178; mov.b32 %r1190, %r1178; mov.b32 %r1191, %r1178; mov.b32 %r1192, %r1178; mov.b32 %r1193, %r1178; mov.b32 %r1194, %r1178; mov.b32 %r1195, %r1178; mov.b32 %r1196, %r1178; mov.b32 %r1197, %r1178; mov.b32 %r1198, %r1178; mov.b32 %r1199, %r1178; mov.b32 %r1200, %r1178; mov.b32 %r1201, %r1178; mov.b32 %r1202, %r1178; mov.b32 %r1203, %r1178; mov.b32 %r1204, %r1178; mov.b32 %r1205, %r1178; mov.b32 %r1206, %r1178; mov.b32 %r1207, %r1178; mov.b32 %r1208, %r1178; mov.b32 %r1209, %r1178; mov.b32 %r1210, %r1178; mov.b32 %r1211, %r1178; mov.b32 %r1212, %r1178; mov.b32 %r1213, %r1178; mov.b32 %r1214, %r1178; mov.b32 %r1215, %r1178; mov.b32 %r1216, %r1178; mov.b32 %r1217, %r1178; mov.b32 %r1218, %r1178; mov.b32 %r1219, %r1178; mov.b32 %r1220, %r1178; mov.b32 %r1221, %r1178; mov.b32 %r1222, %r1178; mov.b32 %r1223, %r1178; mov.b32 %r1224, %r1178; mov.b32 %r1225, %r1178; mov.b32 %r1226, %r1178; mov.b32 %r1227, %r1178; mov.b32 %r1228, %r1178; mov.b32 %r1229, %r1178; mov.b32 %r1230, %r1178; mov.b32 %r1231, %r1178; mov.b32 %r1232, %r1178; mov.b32 %r1233, %r1178; mov.b32 %r1234, %r1178; mov.b32 %r1235, %r1178; mov.b32 %r1236, %r1178; mov.b32 %r1237, %r1178; mov.b32 %r1238, %r1178; mov.b32 %r1239, %r1178; mov.b32 %r1240, %r1178; mov.b32 %r1241, %r1178; mov.b32 %r1243, %r1242; mov.b32 %r1244, %r1242; mov.b32 %r1245, %r1242; mov.b32 %r1174, %r1173; mov.b32 %r1175, %r1173; mov.b32 %r1176, %r1173; $L__BB0_3: // %__nv_exp2f.exit // =>This Inner Loop Header: Depth=1 setp.lt.s32 %p20, %r1172, %r16; add.s32 %r414, %r1170, 1; setp.gt.s32 %p21, %r414, 1; selp.b32 %r1170, 0, %r414, %p21; .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.wait_group 3; bar.sync 0; shl.b32 %r415, %r1170, 12; .loc 1 17 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:17:19 shl.b32 %r416, %r1170, 2; add.s32 %r417, %r126, %r416; ld.shared.b32 %r418, [%r417+32768]; $L__tmp4: .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r160, %r161, %r162, %r163}, [%r17]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r180, %r181, %r182, %r183}, [%r18]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r188, %r189, %r190, %r191}, [%r19]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r240, %r241, %r242, %r243}, [%r20]; $L__tmp5: .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 add.s32 %r419, %r21, %r415; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r164, %r165, %r192, %r193}, [%r419+24576]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r166, %r167, %r198, %r199}, [%r419+25088]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r168, %r169, %r204, %r205}, [%r419+25600]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r170, %r171, %r210, %r211}, [%r419+26112]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r172, %r173, %r216, %r217}, [%r419+26624]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r174, %r175, %r222, %r223}, [%r419+27136]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r176, %r177, %r228, %r229}, [%r419+27648]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r178, %r179, %r234, %r235}, [%r419+28160]; .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mov.b32 %r184, %r159; mov.b32 %r185, %r159; mov.b32 %r186, %r159; mov.b32 %r187, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r184, %r185, %r186, %r187 }, { %r160, %r161, %r162, %r163 }, { %r164, %r165 }, { %r184, %r185, %r186, %r187 }; // end inline asm mov.b32 %r194, %r159; mov.b32 %r195, %r159; mov.b32 %r196, %r159; mov.b32 %r197, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r194, %r195, %r196, %r197 }, { %r160, %r161, %r162, %r163 }, { %r166, %r167 }, { %r194, %r195, %r196, %r197 }; // end inline asm mov.b32 %r200, %r159; mov.b32 %r201, %r159; mov.b32 %r202, %r159; mov.b32 %r203, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r200, %r201, %r202, %r203 }, { %r160, %r161, %r162, %r163 }, { %r168, %r169 }, { %r200, %r201, %r202, %r203 }; // end inline asm mov.b32 %r206, %r159; mov.b32 %r207, %r159; mov.b32 %r208, %r159; mov.b32 %r209, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r206, %r207, %r208, %r209 }, { %r160, %r161, %r162, %r163 }, { %r170, %r171 }, { %r206, %r207, %r208, %r209 }; // end inline asm mov.b32 %r212, %r159; mov.b32 %r213, %r159; mov.b32 %r214, %r159; mov.b32 %r215, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r212, %r213, %r214, %r215 }, { %r160, %r161, %r162, %r163 }, { %r172, %r173 }, { %r212, %r213, %r214, %r215 }; // end inline asm mov.b32 %r218, %r159; mov.b32 %r219, %r159; mov.b32 %r220, %r159; mov.b32 %r221, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r218, %r219, %r220, %r221 }, { %r160, %r161, %r162, %r163 }, { %r174, %r175 }, { %r218, %r219, %r220, %r221 }; // end inline asm mov.b32 %r224, %r159; mov.b32 %r225, %r159; mov.b32 %r226, %r159; mov.b32 %r227, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r224, %r225, %r226, %r227 }, { %r160, %r161, %r162, %r163 }, { %r176, %r177 }, { %r224, %r225, %r226, %r227 }; // end inline asm mov.b32 %r230, %r159; mov.b32 %r231, %r159; mov.b32 %r232, %r159; mov.b32 %r233, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r230, %r231, %r232, %r233 }, { %r160, %r161, %r162, %r163 }, { %r178, %r179 }, { %r230, %r231, %r232, %r233 }; // end inline asm mov.b32 %r236, %r159; mov.b32 %r237, %r159; mov.b32 %r238, %r159; mov.b32 %r239, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r236, %r237, %r238, %r239 }, { %r180, %r181, %r182, %r183 }, { %r164, %r165 }, { %r236, %r237, %r238, %r239 }; // end inline asm mov.b32 %r244, %r159; mov.b32 %r245, %r159; mov.b32 %r246, %r159; mov.b32 %r247, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r244, %r245, %r246, %r247 }, { %r180, %r181, %r182, %r183 }, { %r166, %r167 }, { %r244, %r245, %r246, %r247 }; // end inline asm mov.b32 %r248, %r159; mov.b32 %r249, %r159; mov.b32 %r250, %r159; mov.b32 %r251, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r248, %r249, %r250, %r251 }, { %r180, %r181, %r182, %r183 }, { %r168, %r169 }, { %r248, %r249, %r250, %r251 }; // end inline asm mov.b32 %r252, %r159; mov.b32 %r253, %r159; mov.b32 %r254, %r159; mov.b32 %r255, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r252, %r253, %r254, %r255 }, { %r180, %r181, %r182, %r183 }, { %r170, %r171 }, { %r252, %r253, %r254, %r255 }; // end inline asm mov.b32 %r256, %r159; mov.b32 %r257, %r159; mov.b32 %r258, %r159; mov.b32 %r259, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r256, %r257, %r258, %r259 }, { %r180, %r181, %r182, %r183 }, { %r172, %r173 }, { %r256, %r257, %r258, %r259 }; // end inline asm mov.b32 %r260, %r159; mov.b32 %r261, %r159; mov.b32 %r262, %r159; mov.b32 %r263, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r260, %r261, %r262, %r263 }, { %r180, %r181, %r182, %r183 }, { %r174, %r175 }, { %r260, %r261, %r262, %r263 }; // end inline asm mov.b32 %r264, %r159; mov.b32 %r265, %r159; mov.b32 %r266, %r159; mov.b32 %r267, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r264, %r265, %r266, %r267 }, { %r180, %r181, %r182, %r183 }, { %r176, %r177 }, { %r264, %r265, %r266, %r267 }; // end inline asm mov.b32 %r268, %r159; mov.b32 %r269, %r159; mov.b32 %r270, %r159; mov.b32 %r271, %r159; // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r268, %r269, %r270, %r271 }, { %r180, %r181, %r182, %r183 }, { %r178, %r179 }, { %r268, %r269, %r270, %r271 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r184, %r185, %r186, %r187 }, { %r188, %r189, %r190, %r191 }, { %r192, %r193 }, { %r184, %r185, %r186, %r187 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r194, %r195, %r196, %r197 }, { %r188, %r189, %r190, %r191 }, { %r198, %r199 }, { %r194, %r195, %r196, %r197 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r200, %r201, %r202, %r203 }, { %r188, %r189, %r190, %r191 }, { %r204, %r205 }, { %r200, %r201, %r202, %r203 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r206, %r207, %r208, %r209 }, { %r188, %r189, %r190, %r191 }, { %r210, %r211 }, { %r206, %r207, %r208, %r209 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r212, %r213, %r214, %r215 }, { %r188, %r189, %r190, %r191 }, { %r216, %r217 }, { %r212, %r213, %r214, %r215 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r218, %r219, %r220, %r221 }, { %r188, %r189, %r190, %r191 }, { %r222, %r223 }, { %r218, %r219, %r220, %r221 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r224, %r225, %r226, %r227 }, { %r188, %r189, %r190, %r191 }, { %r228, %r229 }, { %r224, %r225, %r226, %r227 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r230, %r231, %r232, %r233 }, { %r188, %r189, %r190, %r191 }, { %r234, %r235 }, { %r230, %r231, %r232, %r233 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r236, %r237, %r238, %r239 }, { %r240, %r241, %r242, %r243 }, { %r192, %r193 }, { %r236, %r237, %r238, %r239 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r244, %r245, %r246, %r247 }, { %r240, %r241, %r242, %r243 }, { %r198, %r199 }, { %r244, %r245, %r246, %r247 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r248, %r249, %r250, %r251 }, { %r240, %r241, %r242, %r243 }, { %r204, %r205 }, { %r248, %r249, %r250, %r251 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r252, %r253, %r254, %r255 }, { %r240, %r241, %r242, %r243 }, { %r210, %r211 }, { %r252, %r253, %r254, %r255 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r256, %r257, %r258, %r259 }, { %r240, %r241, %r242, %r243 }, { %r216, %r217 }, { %r256, %r257, %r258, %r259 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r260, %r261, %r262, %r263 }, { %r240, %r241, %r242, %r243 }, { %r222, %r223 }, { %r260, %r261, %r262, %r263 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r264, %r265, %r266, %r267 }, { %r240, %r241, %r242, %r243 }, { %r228, %r229 }, { %r264, %r265, %r266, %r267 }; // end inline asm // begin inline asm .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mma.sync.aligned.m16n8k32.row.col.satfinite.s32.s8.s8.s32 { %r268, %r269, %r270, %r271 }, { %r240, %r241, %r242, %r243 }, { %r234, %r235 }, { %r268, %r269, %r270, %r271 }; // end inline asm cvt.rn.f32.s32 %r420, %r231; cvt.rn.f32.s32 %r421, %r230; cvt.rn.f32.s32 %r422, %r225; cvt.rn.f32.s32 %r423, %r224; cvt.rn.f32.s32 %r424, %r219; cvt.rn.f32.s32 %r425, %r218; cvt.rn.f32.s32 %r426, %r213; cvt.rn.f32.s32 %r427, %r212; cvt.rn.f32.s32 %r428, %r207; cvt.rn.f32.s32 %r429, %r206; cvt.rn.f32.s32 %r430, %r201; cvt.rn.f32.s32 %r431, %r200; cvt.rn.f32.s32 %r432, %r195; cvt.rn.f32.s32 %r433, %r194; cvt.rn.f32.s32 %r434, %r185; cvt.rn.f32.s32 %r435, %r184; mul.f32 %r436, %r59, %r435; mul.f32 %r437, %r59, %r434; mul.f32 %r438, %r59, %r433; mul.f32 %r439, %r59, %r432; mul.f32 %r440, %r59, %r431; mul.f32 %r441, %r59, %r430; mul.f32 %r442, %r59, %r429; mul.f32 %r443, %r59, %r428; mul.f32 %r444, %r59, %r427; mul.f32 %r445, %r59, %r426; mul.f32 %r446, %r59, %r425; mul.f32 %r447, %r59, %r424; mul.f32 %r448, %r59, %r423; mul.f32 %r449, %r59, %r422; mul.f32 %r450, %r59, %r421; mul.f32 %r451, %r59, %r420; mul.f32 %r452, %r418, %r451; mul.f32 %r453, %r418, %r450; mul.f32 %r454, %r418, %r449; mul.f32 %r455, %r418, %r448; mul.f32 %r456, %r418, %r447; mul.f32 %r457, %r418, %r446; mul.f32 %r458, %r418, %r445; mul.f32 %r459, %r418, %r444; mul.f32 %r460, %r418, %r443; mul.f32 %r461, %r418, %r442; mul.f32 %r462, %r418, %r441; mul.f32 %r463, %r418, %r440; mul.f32 %r464, %r418, %r439; mul.f32 %r465, %r418, %r438; mul.f32 %r466, %r418, %r437; mul.f32 %r467, %r418, %r436; cvt.rn.f32.s32 %r468, %r233; cvt.rn.f32.s32 %r469, %r232; cvt.rn.f32.s32 %r470, %r227; cvt.rn.f32.s32 %r471, %r226; cvt.rn.f32.s32 %r472, %r221; cvt.rn.f32.s32 %r473, %r220; cvt.rn.f32.s32 %r474, %r215; cvt.rn.f32.s32 %r475, %r214; cvt.rn.f32.s32 %r476, %r209; cvt.rn.f32.s32 %r477, %r208; cvt.rn.f32.s32 %r478, %r203; cvt.rn.f32.s32 %r479, %r202; cvt.rn.f32.s32 %r480, %r197; cvt.rn.f32.s32 %r481, %r196; cvt.rn.f32.s32 %r482, %r187; cvt.rn.f32.s32 %r483, %r186; mul.f32 %r484, %r59, %r483; mul.f32 %r485, %r59, %r482; mul.f32 %r486, %r59, %r481; mul.f32 %r487, %r59, %r480; mul.f32 %r488, %r59, %r479; mul.f32 %r489, %r59, %r478; mul.f32 %r490, %r59, %r477; mul.f32 %r491, %r59, %r476; mul.f32 %r492, %r59, %r475; mul.f32 %r493, %r59, %r474; mul.f32 %r494, %r59, %r473; mul.f32 %r495, %r59, %r472; mul.f32 %r496, %r59, %r471; mul.f32 %r497, %r59, %r470; mul.f32 %r498, %r59, %r469; mul.f32 %r499, %r59, %r468; mul.f32 %r500, %r418, %r499; mul.f32 %r501, %r418, %r498; mul.f32 %r502, %r418, %r497; mul.f32 %r503, %r418, %r496; mul.f32 %r504, %r418, %r495; mul.f32 %r505, %r418, %r494; mul.f32 %r506, %r418, %r493; mul.f32 %r507, %r418, %r492; mul.f32 %r508, %r418, %r491; mul.f32 %r509, %r418, %r490; mul.f32 %r510, %r418, %r489; mul.f32 %r511, %r418, %r488; mul.f32 %r512, %r418, %r487; mul.f32 %r513, %r418, %r486; mul.f32 %r514, %r418, %r485; mul.f32 %r515, %r418, %r484; cvt.rn.f32.s32 %r516, %r269; cvt.rn.f32.s32 %r517, %r268; cvt.rn.f32.s32 %r518, %r265; cvt.rn.f32.s32 %r519, %r264; cvt.rn.f32.s32 %r520, %r261; cvt.rn.f32.s32 %r521, %r260; cvt.rn.f32.s32 %r522, %r257; cvt.rn.f32.s32 %r523, %r256; cvt.rn.f32.s32 %r524, %r253; cvt.rn.f32.s32 %r525, %r252; cvt.rn.f32.s32 %r526, %r249; cvt.rn.f32.s32 %r527, %r248; cvt.rn.f32.s32 %r528, %r245; cvt.rn.f32.s32 %r529, %r244; cvt.rn.f32.s32 %r530, %r237; cvt.rn.f32.s32 %r531, %r236; mul.f32 %r532, %r59, %r531; mul.f32 %r533, %r59, %r530; mul.f32 %r534, %r59, %r529; mul.f32 %r535, %r59, %r528; mul.f32 %r536, %r59, %r527; mul.f32 %r537, %r59, %r526; mul.f32 %r538, %r59, %r525; mul.f32 %r539, %r59, %r524; mul.f32 %r540, %r59, %r523; mul.f32 %r541, %r59, %r522; mul.f32 %r542, %r59, %r521; mul.f32 %r543, %r59, %r520; mul.f32 %r544, %r59, %r519; mul.f32 %r545, %r59, %r518; mul.f32 %r546, %r59, %r517; mul.f32 %r547, %r59, %r516; mul.f32 %r548, %r418, %r547; mul.f32 %r549, %r418, %r546; mul.f32 %r550, %r418, %r545; mul.f32 %r551, %r418, %r544; mul.f32 %r552, %r418, %r543; mul.f32 %r553, %r418, %r542; mul.f32 %r554, %r418, %r541; mul.f32 %r555, %r418, %r540; mul.f32 %r556, %r418, %r539; mul.f32 %r557, %r418, %r538; mul.f32 %r558, %r418, %r537; mul.f32 %r559, %r418, %r536; mul.f32 %r560, %r418, %r535; mul.f32 %r561, %r418, %r534; mul.f32 %r562, %r418, %r533; mul.f32 %r563, %r418, %r532; cvt.rn.f32.s32 %r564, %r271; cvt.rn.f32.s32 %r565, %r270; cvt.rn.f32.s32 %r566, %r267; cvt.rn.f32.s32 %r567, %r266; cvt.rn.f32.s32 %r568, %r263; cvt.rn.f32.s32 %r569, %r262; cvt.rn.f32.s32 %r570, %r259; cvt.rn.f32.s32 %r571, %r258; cvt.rn.f32.s32 %r572, %r255; cvt.rn.f32.s32 %r573, %r254; cvt.rn.f32.s32 %r574, %r251; cvt.rn.f32.s32 %r575, %r250; cvt.rn.f32.s32 %r576, %r247; cvt.rn.f32.s32 %r577, %r246; cvt.rn.f32.s32 %r578, %r239; cvt.rn.f32.s32 %r579, %r238; mul.f32 %r580, %r59, %r579; mul.f32 %r581, %r59, %r578; mul.f32 %r582, %r59, %r577; mul.f32 %r583, %r59, %r576; mul.f32 %r584, %r59, %r575; mul.f32 %r585, %r59, %r574; mul.f32 %r586, %r59, %r573; mul.f32 %r587, %r59, %r572; mul.f32 %r588, %r59, %r571; mul.f32 %r589, %r59, %r570; mul.f32 %r590, %r59, %r569; mul.f32 %r591, %r59, %r568; mul.f32 %r592, %r59, %r567; mul.f32 %r593, %r59, %r566; mul.f32 %r594, %r59, %r565; mul.f32 %r595, %r59, %r564; mul.f32 %r596, %r418, %r595; mul.f32 %r597, %r418, %r594; mul.f32 %r598, %r418, %r593; mul.f32 %r599, %r418, %r592; mul.f32 %r600, %r418, %r591; mul.f32 %r601, %r418, %r590; mul.f32 %r602, %r418, %r589; mul.f32 %r603, %r418, %r588; mul.f32 %r604, %r418, %r587; mul.f32 %r605, %r418, %r586; mul.f32 %r606, %r418, %r585; mul.f32 %r607, %r418, %r584; mul.f32 %r608, %r418, %r583; mul.f32 %r609, %r418, %r582; mul.f32 %r610, %r418, %r581; mul.f32 %r611, %r418, %r580; $L__tmp6: .loc 1 19 32, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:19:32 .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r612, %r467, %r466; max.f32 %r613, %r465, %r464; max.f32 %r614, %r463, %r462; max.f32 %r615, %r461, %r460; max.f32 %r616, %r459, %r458; max.f32 %r617, %r457, %r456; max.f32 %r618, %r455, %r454; max.f32 %r619, %r453, %r452; max.f32 %r620, %r612, %r613; max.f32 %r621, %r614, %r615; max.f32 %r622, %r616, %r617; max.f32 %r623, %r618, %r619; max.f32 %r624, %r620, %r621; max.f32 %r625, %r622, %r623; max.f32 %r626, %r624, %r625; max.f32 %r627, %r515, %r514; max.f32 %r628, %r513, %r512; max.f32 %r629, %r511, %r510; max.f32 %r630, %r509, %r508; max.f32 %r631, %r507, %r506; max.f32 %r632, %r505, %r504; max.f32 %r633, %r503, %r502; max.f32 %r634, %r501, %r500; max.f32 %r635, %r627, %r628; max.f32 %r636, %r629, %r630; max.f32 %r637, %r631, %r632; max.f32 %r638, %r633, %r634; max.f32 %r639, %r635, %r636; max.f32 %r640, %r637, %r638; max.f32 %r641, %r639, %r640; max.f32 %r642, %r563, %r562; max.f32 %r643, %r561, %r560; max.f32 %r644, %r559, %r558; max.f32 %r645, %r557, %r556; max.f32 %r646, %r555, %r554; max.f32 %r647, %r553, %r552; max.f32 %r648, %r551, %r550; max.f32 %r649, %r549, %r548; max.f32 %r650, %r642, %r643; max.f32 %r651, %r644, %r645; max.f32 %r652, %r646, %r647; max.f32 %r653, %r648, %r649; max.f32 %r654, %r650, %r651; max.f32 %r655, %r652, %r653; max.f32 %r656, %r654, %r655; max.f32 %r657, %r611, %r610; max.f32 %r658, %r609, %r608; max.f32 %r659, %r607, %r606; max.f32 %r660, %r605, %r604; max.f32 %r661, %r603, %r602; max.f32 %r662, %r601, %r600; max.f32 %r663, %r599, %r598; max.f32 %r664, %r597, %r596; max.f32 %r665, %r657, %r658; max.f32 %r666, %r659, %r660; max.f32 %r667, %r661, %r662; max.f32 %r668, %r663, %r664; max.f32 %r669, %r665, %r666; max.f32 %r670, %r667, %r668; max.f32 %r671, %r669, %r670; $L__tmp7: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r672, %r626, 2, 31, -1; $L__tmp8: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r673, %r626, %r672; $L__tmp9: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r674, %r673, 1, 31, -1; shfl.sync.bfly.b32 %r675, %r641, 2, 31, -1; $L__tmp10: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r676, %r641, %r675; $L__tmp11: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r677, %r676, 1, 31, -1; shfl.sync.bfly.b32 %r678, %r656, 2, 31, -1; $L__tmp12: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r679, %r656, %r678; $L__tmp13: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r680, %r679, 1, 31, -1; shfl.sync.bfly.b32 %r681, %r671, 2, 31, -1; $L__tmp14: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r682, %r671, %r681; $L__tmp15: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r683, %r682, 1, 31, -1; $L__tmp16: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r684, %r673, %r674; max.f32 %r685, %r676, %r677; max.f32 %r686, %r679, %r680; max.f32 %r687, %r682, %r683; $L__tmp17: .loc 1 19 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:19:16 max.f32 %r33, %r1176, %r687; max.f32 %r32, %r1175, %r686; max.f32 %r31, %r1174, %r685; max.f32 %r30, %r1173, %r684; .loc 1 20 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:20:14 neg.f32 %r688, %r30; fma.rn.f32 %r689, %r418, %r436, %r688; fma.rn.f32 %r690, %r418, %r437, %r688; neg.f32 %r691, %r31; fma.rn.f32 %r692, %r418, %r484, %r691; fma.rn.f32 %r693, %r418, %r485, %r691; fma.rn.f32 %r694, %r418, %r438, %r688; fma.rn.f32 %r695, %r418, %r439, %r688; fma.rn.f32 %r696, %r418, %r486, %r691; fma.rn.f32 %r697, %r418, %r487, %r691; fma.rn.f32 %r698, %r418, %r440, %r688; fma.rn.f32 %r699, %r418, %r441, %r688; fma.rn.f32 %r700, %r418, %r488, %r691; fma.rn.f32 %r701, %r418, %r489, %r691; fma.rn.f32 %r702, %r418, %r442, %r688; fma.rn.f32 %r703, %r418, %r443, %r688; fma.rn.f32 %r704, %r418, %r490, %r691; fma.rn.f32 %r705, %r418, %r491, %r691; fma.rn.f32 %r706, %r418, %r444, %r688; fma.rn.f32 %r707, %r418, %r445, %r688; fma.rn.f32 %r708, %r418, %r492, %r691; fma.rn.f32 %r709, %r418, %r493, %r691; fma.rn.f32 %r710, %r418, %r446, %r688; fma.rn.f32 %r711, %r418, %r447, %r688; fma.rn.f32 %r712, %r418, %r494, %r691; fma.rn.f32 %r713, %r418, %r495, %r691; fma.rn.f32 %r714, %r418, %r448, %r688; fma.rn.f32 %r715, %r418, %r449, %r688; fma.rn.f32 %r716, %r418, %r496, %r691; fma.rn.f32 %r717, %r418, %r497, %r691; fma.rn.f32 %r718, %r418, %r450, %r688; fma.rn.f32 %r719, %r418, %r451, %r688; fma.rn.f32 %r720, %r418, %r498, %r691; fma.rn.f32 %r721, %r418, %r499, %r691; neg.f32 %r722, %r32; fma.rn.f32 %r723, %r418, %r532, %r722; fma.rn.f32 %r724, %r418, %r533, %r722; neg.f32 %r725, %r33; fma.rn.f32 %r726, %r418, %r580, %r725; fma.rn.f32 %r727, %r418, %r581, %r725; fma.rn.f32 %r728, %r418, %r534, %r722; fma.rn.f32 %r729, %r418, %r535, %r722; fma.rn.f32 %r730, %r418, %r582, %r725; fma.rn.f32 %r731, %r418, %r583, %r725; fma.rn.f32 %r732, %r418, %r536, %r722; fma.rn.f32 %r733, %r418, %r537, %r722; fma.rn.f32 %r734, %r418, %r584, %r725; fma.rn.f32 %r735, %r418, %r585, %r725; fma.rn.f32 %r736, %r418, %r538, %r722; fma.rn.f32 %r737, %r418, %r539, %r722; fma.rn.f32 %r738, %r418, %r586, %r725; fma.rn.f32 %r739, %r418, %r587, %r725; fma.rn.f32 %r740, %r418, %r540, %r722; fma.rn.f32 %r741, %r418, %r541, %r722; fma.rn.f32 %r742, %r418, %r588, %r725; fma.rn.f32 %r743, %r418, %r589, %r725; fma.rn.f32 %r744, %r418, %r542, %r722; fma.rn.f32 %r745, %r418, %r543, %r722; fma.rn.f32 %r746, %r418, %r590, %r725; fma.rn.f32 %r747, %r418, %r591, %r725; fma.rn.f32 %r748, %r418, %r544, %r722; fma.rn.f32 %r749, %r418, %r545, %r722; fma.rn.f32 %r750, %r418, %r592, %r725; fma.rn.f32 %r751, %r418, %r593, %r725; fma.rn.f32 %r752, %r418, %r546, %r722; fma.rn.f32 %r753, %r418, %r547, %r722; fma.rn.f32 %r754, %r418, %r594, %r725; fma.rn.f32 %r755, %r418, %r595, %r725; .loc 1 21 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:21:13 ex2.approx.ftz.f32 %r756, %r689; ex2.approx.ftz.f32 %r757, %r690; ex2.approx.ftz.f32 %r758, %r692; ex2.approx.ftz.f32 %r759, %r693; ex2.approx.ftz.f32 %r760, %r694; ex2.approx.ftz.f32 %r761, %r695; ex2.approx.ftz.f32 %r762, %r696; ex2.approx.ftz.f32 %r763, %r697; ex2.approx.ftz.f32 %r764, %r698; ex2.approx.ftz.f32 %r765, %r699; ex2.approx.ftz.f32 %r766, %r700; ex2.approx.ftz.f32 %r767, %r701; ex2.approx.ftz.f32 %r768, %r702; ex2.approx.ftz.f32 %r769, %r703; ex2.approx.ftz.f32 %r770, %r704; ex2.approx.ftz.f32 %r771, %r705; ex2.approx.ftz.f32 %r772, %r706; ex2.approx.ftz.f32 %r773, %r707; ex2.approx.ftz.f32 %r774, %r708; ex2.approx.ftz.f32 %r775, %r709; ex2.approx.ftz.f32 %r776, %r710; ex2.approx.ftz.f32 %r777, %r711; ex2.approx.ftz.f32 %r778, %r712; ex2.approx.ftz.f32 %r779, %r713; ex2.approx.ftz.f32 %r780, %r714; ex2.approx.ftz.f32 %r781, %r715; ex2.approx.ftz.f32 %r782, %r716; ex2.approx.ftz.f32 %r783, %r717; ex2.approx.ftz.f32 %r784, %r718; ex2.approx.ftz.f32 %r785, %r719; ex2.approx.ftz.f32 %r786, %r720; ex2.approx.ftz.f32 %r787, %r721; ex2.approx.ftz.f32 %r788, %r723; ex2.approx.ftz.f32 %r789, %r724; ex2.approx.ftz.f32 %r790, %r726; ex2.approx.ftz.f32 %r791, %r727; ex2.approx.ftz.f32 %r792, %r728; ex2.approx.ftz.f32 %r793, %r729; ex2.approx.ftz.f32 %r794, %r730; ex2.approx.ftz.f32 %r795, %r731; ex2.approx.ftz.f32 %r796, %r732; ex2.approx.ftz.f32 %r797, %r733; ex2.approx.ftz.f32 %r798, %r734; ex2.approx.ftz.f32 %r799, %r735; ex2.approx.ftz.f32 %r800, %r736; ex2.approx.ftz.f32 %r801, %r737; ex2.approx.ftz.f32 %r802, %r738; ex2.approx.ftz.f32 %r803, %r739; ex2.approx.ftz.f32 %r804, %r740; ex2.approx.ftz.f32 %r805, %r741; ex2.approx.ftz.f32 %r806, %r742; ex2.approx.ftz.f32 %r807, %r743; ex2.approx.ftz.f32 %r808, %r744; ex2.approx.ftz.f32 %r809, %r745; ex2.approx.ftz.f32 %r810, %r746; ex2.approx.ftz.f32 %r811, %r747; ex2.approx.ftz.f32 %r812, %r748; ex2.approx.ftz.f32 %r813, %r749; ex2.approx.ftz.f32 %r814, %r750; ex2.approx.ftz.f32 %r815, %r751; ex2.approx.ftz.f32 %r816, %r752; ex2.approx.ftz.f32 %r817, %r753; ex2.approx.ftz.f32 %r818, %r754; ex2.approx.ftz.f32 %r819, %r755; $L__tmp18: .loc 1 22 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:22:16 .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r820, %r756, %r757; add.f32 %r821, %r760, %r761; add.f32 %r822, %r764, %r765; add.f32 %r823, %r768, %r769; add.f32 %r824, %r772, %r773; add.f32 %r825, %r776, %r777; add.f32 %r826, %r780, %r781; add.f32 %r827, %r784, %r785; add.f32 %r828, %r820, %r821; add.f32 %r829, %r822, %r823; add.f32 %r830, %r824, %r825; add.f32 %r831, %r826, %r827; add.f32 %r832, %r828, %r829; add.f32 %r833, %r830, %r831; add.f32 %r834, %r832, %r833; add.f32 %r835, %r758, %r759; add.f32 %r836, %r762, %r763; add.f32 %r837, %r766, %r767; add.f32 %r838, %r770, %r771; add.f32 %r839, %r774, %r775; add.f32 %r840, %r778, %r779; add.f32 %r841, %r782, %r783; add.f32 %r842, %r786, %r787; add.f32 %r843, %r835, %r836; add.f32 %r844, %r837, %r838; add.f32 %r845, %r839, %r840; add.f32 %r846, %r841, %r842; add.f32 %r847, %r843, %r844; add.f32 %r848, %r845, %r846; add.f32 %r849, %r847, %r848; add.f32 %r850, %r788, %r789; add.f32 %r851, %r792, %r793; add.f32 %r852, %r796, %r797; add.f32 %r853, %r800, %r801; add.f32 %r854, %r804, %r805; add.f32 %r855, %r808, %r809; add.f32 %r856, %r812, %r813; add.f32 %r857, %r816, %r817; add.f32 %r858, %r850, %r851; add.f32 %r859, %r852, %r853; add.f32 %r860, %r854, %r855; add.f32 %r861, %r856, %r857; add.f32 %r862, %r858, %r859; add.f32 %r863, %r860, %r861; add.f32 %r864, %r862, %r863; add.f32 %r865, %r790, %r791; add.f32 %r866, %r794, %r795; add.f32 %r867, %r798, %r799; add.f32 %r868, %r802, %r803; add.f32 %r869, %r806, %r807; add.f32 %r870, %r810, %r811; add.f32 %r871, %r814, %r815; add.f32 %r872, %r818, %r819; add.f32 %r873, %r865, %r866; add.f32 %r874, %r867, %r868; add.f32 %r875, %r869, %r870; add.f32 %r876, %r871, %r872; add.f32 %r877, %r873, %r874; add.f32 %r878, %r875, %r876; add.f32 %r879, %r877, %r878; $L__tmp19: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r880, %r834, 2, 31, -1; $L__tmp20: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r881, %r834, %r880; $L__tmp21: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r882, %r881, 1, 31, -1; shfl.sync.bfly.b32 %r883, %r849, 2, 31, -1; $L__tmp22: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r884, %r849, %r883; $L__tmp23: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r885, %r884, 1, 31, -1; shfl.sync.bfly.b32 %r886, %r864, 2, 31, -1; $L__tmp24: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r887, %r864, %r886; $L__tmp25: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r888, %r887, 1, 31, -1; shfl.sync.bfly.b32 %r889, %r879, 2, 31, -1; $L__tmp26: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r890, %r879, %r889; $L__tmp27: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r891, %r890, 1, 31, -1; $L__tmp28: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r892, %r890, %r891; add.f32 %r893, %r887, %r888; add.f32 %r894, %r884, %r885; add.f32 %r895, %r881, %r882; $L__tmp29: .loc 1 24 30, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:24:30 sub.f32 %r896, %r1176, %r33; sub.f32 %r897, %r1175, %r32; sub.f32 %r898, %r1174, %r31; sub.f32 %r899, %r1173, %r30; .loc 1 24 17, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:24:17 ex2.approx.ftz.f32 %r900, %r899; ex2.approx.ftz.f32 %r901, %r898; ex2.approx.ftz.f32 %r902, %r897; ex2.approx.ftz.f32 %r903, %r896; .loc 1 25 15, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:25:15 fma.rn.f32 %r1242, %r1242, %r900, %r895; fma.rn.f32 %r1243, %r1243, %r901, %r894; fma.rn.f32 %r1244, %r1244, %r902, %r893; fma.rn.f32 %r1245, %r1245, %r903, %r892; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 shl.b32 %r904, %r1170, 13; add.s32 %r905, %r126, %r904; add.s32 %r906, %r905, %r22; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r272, %r273, %r298, %r299}, [%r906]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r352, %r353, %r376, %r377}, [%r906+4096]; add.s32 %r907, %r905, %r23; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r278, %r279, %r306, %r307}, [%r907]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r358, %r359, %r382, %r383}, [%r907+4096]; add.s32 %r908, %r905, %r24; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r280, %r281, %r310, %r311}, [%r908]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r360, %r361, %r384, %r385}, [%r908+4096]; add.s32 %r909, %r905, %r25; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r282, %r283, %r314, %r315}, [%r909]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r362, %r363, %r386, %r387}, [%r909+4096]; add.s32 %r910, %r905, %r26; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r284, %r285, %r318, %r319}, [%r910]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r364, %r365, %r388, %r389}, [%r910+4096]; add.s32 %r911, %r905, %r27; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r286, %r287, %r322, %r323}, [%r911]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r366, %r367, %r390, %r391}, [%r911+4096]; add.s32 %r912, %r905, %r28; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r288, %r289, %r326, %r327}, [%r912]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r368, %r369, %r392, %r393}, [%r912+4096]; add.s32 %r913, %r905, %r29; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r290, %r291, %r330, %r331}, [%r913]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r370, %r371, %r394, %r395}, [%r913+4096]; .loc 1 30 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:30:13 cvt.rn.f16x2.f32 %r274, %r757, %r756; cvt.rn.f16x2.f32 %r275, %r759, %r758; cvt.rn.f16x2.f32 %r276, %r761, %r760; cvt.rn.f16x2.f32 %r277, %r763, %r762; cvt.rn.f16x2.f32 %r302, %r765, %r764; cvt.rn.f16x2.f32 %r303, %r767, %r766; cvt.rn.f16x2.f32 %r304, %r769, %r768; cvt.rn.f16x2.f32 %r305, %r771, %r770; cvt.rn.f16x2.f32 %r354, %r773, %r772; cvt.rn.f16x2.f32 %r355, %r775, %r774; cvt.rn.f16x2.f32 %r356, %r777, %r776; cvt.rn.f16x2.f32 %r357, %r779, %r778; cvt.rn.f16x2.f32 %r378, %r781, %r780; cvt.rn.f16x2.f32 %r379, %r783, %r782; cvt.rn.f16x2.f32 %r380, %r785, %r784; cvt.rn.f16x2.f32 %r381, %r787, %r786; cvt.rn.f16x2.f32 %r292, %r789, %r788; cvt.rn.f16x2.f32 %r293, %r791, %r790; cvt.rn.f16x2.f32 %r294, %r793, %r792; cvt.rn.f16x2.f32 %r295, %r795, %r794; cvt.rn.f16x2.f32 %r336, %r797, %r796; cvt.rn.f16x2.f32 %r337, %r799, %r798; cvt.rn.f16x2.f32 %r338, %r801, %r800; cvt.rn.f16x2.f32 %r339, %r803, %r802; cvt.rn.f16x2.f32 %r372, %r805, %r804; cvt.rn.f16x2.f32 %r373, %r807, %r806; cvt.rn.f16x2.f32 %r374, %r809, %r808; cvt.rn.f16x2.f32 %r375, %r811, %r810; cvt.rn.f16x2.f32 %r396, %r813, %r812; cvt.rn.f16x2.f32 %r397, %r815, %r814; cvt.rn.f16x2.f32 %r398, %r817, %r816; cvt.rn.f16x2.f32 %r399, %r819, %r818; .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mov.b32 %r296, %r159; mov.b32 %r297, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r296, %r297 }, { %r274, %r275, %r276, %r277 }, { %r272, %r273 }, { %r296, %r297 }; // end inline asm mov.b32 %r300, %r159; mov.b32 %r301, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r300, %r301 }, { %r274, %r275, %r276, %r277 }, { %r278, %r279 }, { %r300, %r301 }; // end inline asm mov.b32 %r308, %r159; mov.b32 %r309, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r308, %r309 }, { %r274, %r275, %r276, %r277 }, { %r280, %r281 }, { %r308, %r309 }; // end inline asm mov.b32 %r312, %r159; mov.b32 %r313, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r312, %r313 }, { %r274, %r275, %r276, %r277 }, { %r282, %r283 }, { %r312, %r313 }; // end inline asm mov.b32 %r316, %r159; mov.b32 %r317, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r316, %r317 }, { %r274, %r275, %r276, %r277 }, { %r284, %r285 }, { %r316, %r317 }; // end inline asm mov.b32 %r320, %r159; mov.b32 %r321, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r320, %r321 }, { %r274, %r275, %r276, %r277 }, { %r286, %r287 }, { %r320, %r321 }; // end inline asm mov.b32 %r324, %r159; mov.b32 %r325, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r324, %r325 }, { %r274, %r275, %r276, %r277 }, { %r288, %r289 }, { %r324, %r325 }; // end inline asm mov.b32 %r328, %r159; mov.b32 %r329, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r328, %r329 }, { %r274, %r275, %r276, %r277 }, { %r290, %r291 }, { %r328, %r329 }; // end inline asm mov.b32 %r332, %r159; mov.b32 %r333, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r332, %r333 }, { %r292, %r293, %r294, %r295 }, { %r272, %r273 }, { %r332, %r333 }; // end inline asm mov.b32 %r334, %r159; mov.b32 %r335, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r334, %r335 }, { %r292, %r293, %r294, %r295 }, { %r278, %r279 }, { %r334, %r335 }; // end inline asm mov.b32 %r340, %r159; mov.b32 %r341, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r340, %r341 }, { %r292, %r293, %r294, %r295 }, { %r280, %r281 }, { %r340, %r341 }; // end inline asm mov.b32 %r342, %r159; mov.b32 %r343, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r342, %r343 }, { %r292, %r293, %r294, %r295 }, { %r282, %r283 }, { %r342, %r343 }; // end inline asm mov.b32 %r344, %r159; mov.b32 %r345, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r344, %r345 }, { %r292, %r293, %r294, %r295 }, { %r284, %r285 }, { %r344, %r345 }; // end inline asm mov.b32 %r346, %r159; mov.b32 %r347, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r346, %r347 }, { %r292, %r293, %r294, %r295 }, { %r286, %r287 }, { %r346, %r347 }; // end inline asm mov.b32 %r348, %r159; mov.b32 %r349, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r348, %r349 }, { %r292, %r293, %r294, %r295 }, { %r288, %r289 }, { %r348, %r349 }; // end inline asm mov.b32 %r350, %r159; mov.b32 %r351, %r159; // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r350, %r351 }, { %r292, %r293, %r294, %r295 }, { %r290, %r291 }, { %r350, %r351 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r296, %r297 }, { %r302, %r303, %r304, %r305 }, { %r298, %r299 }, { %r296, %r297 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r300, %r301 }, { %r302, %r303, %r304, %r305 }, { %r306, %r307 }, { %r300, %r301 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r308, %r309 }, { %r302, %r303, %r304, %r305 }, { %r310, %r311 }, { %r308, %r309 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r312, %r313 }, { %r302, %r303, %r304, %r305 }, { %r314, %r315 }, { %r312, %r313 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r316, %r317 }, { %r302, %r303, %r304, %r305 }, { %r318, %r319 }, { %r316, %r317 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r320, %r321 }, { %r302, %r303, %r304, %r305 }, { %r322, %r323 }, { %r320, %r321 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r324, %r325 }, { %r302, %r303, %r304, %r305 }, { %r326, %r327 }, { %r324, %r325 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r328, %r329 }, { %r302, %r303, %r304, %r305 }, { %r330, %r331 }, { %r328, %r329 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r332, %r333 }, { %r336, %r337, %r338, %r339 }, { %r298, %r299 }, { %r332, %r333 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r334, %r335 }, { %r336, %r337, %r338, %r339 }, { %r306, %r307 }, { %r334, %r335 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r340, %r341 }, { %r336, %r337, %r338, %r339 }, { %r310, %r311 }, { %r340, %r341 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r342, %r343 }, { %r336, %r337, %r338, %r339 }, { %r314, %r315 }, { %r342, %r343 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r344, %r345 }, { %r336, %r337, %r338, %r339 }, { %r318, %r319 }, { %r344, %r345 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r346, %r347 }, { %r336, %r337, %r338, %r339 }, { %r322, %r323 }, { %r346, %r347 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r348, %r349 }, { %r336, %r337, %r338, %r339 }, { %r326, %r327 }, { %r348, %r349 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r350, %r351 }, { %r336, %r337, %r338, %r339 }, { %r330, %r331 }, { %r350, %r351 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r296, %r297 }, { %r354, %r355, %r356, %r357 }, { %r352, %r353 }, { %r296, %r297 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r300, %r301 }, { %r354, %r355, %r356, %r357 }, { %r358, %r359 }, { %r300, %r301 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r308, %r309 }, { %r354, %r355, %r356, %r357 }, { %r360, %r361 }, { %r308, %r309 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r312, %r313 }, { %r354, %r355, %r356, %r357 }, { %r362, %r363 }, { %r312, %r313 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r316, %r317 }, { %r354, %r355, %r356, %r357 }, { %r364, %r365 }, { %r316, %r317 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r320, %r321 }, { %r354, %r355, %r356, %r357 }, { %r366, %r367 }, { %r320, %r321 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r324, %r325 }, { %r354, %r355, %r356, %r357 }, { %r368, %r369 }, { %r324, %r325 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r328, %r329 }, { %r354, %r355, %r356, %r357 }, { %r370, %r371 }, { %r328, %r329 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r332, %r333 }, { %r372, %r373, %r374, %r375 }, { %r352, %r353 }, { %r332, %r333 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r334, %r335 }, { %r372, %r373, %r374, %r375 }, { %r358, %r359 }, { %r334, %r335 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r340, %r341 }, { %r372, %r373, %r374, %r375 }, { %r360, %r361 }, { %r340, %r341 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r342, %r343 }, { %r372, %r373, %r374, %r375 }, { %r362, %r363 }, { %r342, %r343 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r344, %r345 }, { %r372, %r373, %r374, %r375 }, { %r364, %r365 }, { %r344, %r345 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r346, %r347 }, { %r372, %r373, %r374, %r375 }, { %r366, %r367 }, { %r346, %r347 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r348, %r349 }, { %r372, %r373, %r374, %r375 }, { %r368, %r369 }, { %r348, %r349 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r350, %r351 }, { %r372, %r373, %r374, %r375 }, { %r370, %r371 }, { %r350, %r351 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r296, %r297 }, { %r378, %r379, %r380, %r381 }, { %r376, %r377 }, { %r296, %r297 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r300, %r301 }, { %r378, %r379, %r380, %r381 }, { %r382, %r383 }, { %r300, %r301 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r308, %r309 }, { %r378, %r379, %r380, %r381 }, { %r384, %r385 }, { %r308, %r309 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r312, %r313 }, { %r378, %r379, %r380, %r381 }, { %r386, %r387 }, { %r312, %r313 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r316, %r317 }, { %r378, %r379, %r380, %r381 }, { %r388, %r389 }, { %r316, %r317 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r320, %r321 }, { %r378, %r379, %r380, %r381 }, { %r390, %r391 }, { %r320, %r321 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r324, %r325 }, { %r378, %r379, %r380, %r381 }, { %r392, %r393 }, { %r324, %r325 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r328, %r329 }, { %r378, %r379, %r380, %r381 }, { %r394, %r395 }, { %r328, %r329 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r332, %r333 }, { %r396, %r397, %r398, %r399 }, { %r376, %r377 }, { %r332, %r333 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r334, %r335 }, { %r396, %r397, %r398, %r399 }, { %r382, %r383 }, { %r334, %r335 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r340, %r341 }, { %r396, %r397, %r398, %r399 }, { %r384, %r385 }, { %r340, %r341 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r342, %r343 }, { %r396, %r397, %r398, %r399 }, { %r386, %r387 }, { %r342, %r343 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r344, %r345 }, { %r396, %r397, %r398, %r399 }, { %r388, %r389 }, { %r344, %r345 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r346, %r347 }, { %r396, %r397, %r398, %r399 }, { %r390, %r391 }, { %r346, %r347 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r348, %r349 }, { %r396, %r397, %r398, %r399 }, { %r392, %r393 }, { %r348, %r349 }; // end inline asm // begin inline asm .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 { %r350, %r351 }, { %r396, %r397, %r398, %r399 }, { %r394, %r395 }, { %r350, %r351 }; // end inline asm .loc 1 32 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:9 mov.b32 {%rs1, %rs2}, %r350; cvt.f32.f16 %r914, %rs1; cvt.f32.f16 %r915, %rs2; mov.b32 {%rs3, %rs4}, %r351; cvt.f32.f16 %r916, %rs3; cvt.f32.f16 %r917, %rs4; mov.b32 {%rs5, %rs6}, %r348; cvt.f32.f16 %r918, %rs5; cvt.f32.f16 %r919, %rs6; mov.b32 {%rs7, %rs8}, %r349; cvt.f32.f16 %r920, %rs7; cvt.f32.f16 %r921, %rs8; mov.b32 {%rs9, %rs10}, %r346; cvt.f32.f16 %r922, %rs9; cvt.f32.f16 %r923, %rs10; mov.b32 {%rs11, %rs12}, %r347; cvt.f32.f16 %r924, %rs11; cvt.f32.f16 %r925, %rs12; mov.b32 {%rs13, %rs14}, %r344; cvt.f32.f16 %r926, %rs13; cvt.f32.f16 %r927, %rs14; mov.b32 {%rs15, %rs16}, %r345; cvt.f32.f16 %r928, %rs15; cvt.f32.f16 %r929, %rs16; mov.b32 {%rs17, %rs18}, %r342; cvt.f32.f16 %r930, %rs17; cvt.f32.f16 %r931, %rs18; mov.b32 {%rs19, %rs20}, %r343; cvt.f32.f16 %r932, %rs19; cvt.f32.f16 %r933, %rs20; mov.b32 {%rs21, %rs22}, %r340; cvt.f32.f16 %r934, %rs21; cvt.f32.f16 %r935, %rs22; mov.b32 {%rs23, %rs24}, %r341; cvt.f32.f16 %r936, %rs23; cvt.f32.f16 %r937, %rs24; mov.b32 {%rs25, %rs26}, %r334; cvt.f32.f16 %r938, %rs25; cvt.f32.f16 %r939, %rs26; mov.b32 {%rs27, %rs28}, %r335; cvt.f32.f16 %r940, %rs27; cvt.f32.f16 %r941, %rs28; mov.b32 {%rs29, %rs30}, %r332; cvt.f32.f16 %r942, %rs29; cvt.f32.f16 %r943, %rs30; mov.b32 {%rs31, %rs32}, %r333; cvt.f32.f16 %r944, %rs31; cvt.f32.f16 %r945, %rs32; mov.b32 {%rs33, %rs34}, %r328; cvt.f32.f16 %r946, %rs33; cvt.f32.f16 %r947, %rs34; mov.b32 {%rs35, %rs36}, %r329; cvt.f32.f16 %r948, %rs35; cvt.f32.f16 %r949, %rs36; mov.b32 {%rs37, %rs38}, %r324; cvt.f32.f16 %r950, %rs37; cvt.f32.f16 %r951, %rs38; mov.b32 {%rs39, %rs40}, %r325; cvt.f32.f16 %r952, %rs39; cvt.f32.f16 %r953, %rs40; mov.b32 {%rs41, %rs42}, %r320; cvt.f32.f16 %r954, %rs41; cvt.f32.f16 %r955, %rs42; mov.b32 {%rs43, %rs44}, %r321; cvt.f32.f16 %r956, %rs43; cvt.f32.f16 %r957, %rs44; mov.b32 {%rs45, %rs46}, %r316; cvt.f32.f16 %r958, %rs45; cvt.f32.f16 %r959, %rs46; mov.b32 {%rs47, %rs48}, %r317; cvt.f32.f16 %r960, %rs47; cvt.f32.f16 %r961, %rs48; mov.b32 {%rs49, %rs50}, %r312; cvt.f32.f16 %r962, %rs49; cvt.f32.f16 %r963, %rs50; mov.b32 {%rs51, %rs52}, %r313; cvt.f32.f16 %r964, %rs51; cvt.f32.f16 %r965, %rs52; mov.b32 {%rs53, %rs54}, %r308; cvt.f32.f16 %r966, %rs53; cvt.f32.f16 %r967, %rs54; mov.b32 {%rs55, %rs56}, %r309; cvt.f32.f16 %r968, %rs55; cvt.f32.f16 %r969, %rs56; mov.b32 {%rs57, %rs58}, %r300; cvt.f32.f16 %r970, %rs57; cvt.f32.f16 %r971, %rs58; mov.b32 {%rs59, %rs60}, %r301; cvt.f32.f16 %r972, %rs59; cvt.f32.f16 %r973, %rs60; mov.b32 {%rs61, %rs62}, %r296; cvt.f32.f16 %r974, %rs61; cvt.f32.f16 %r975, %rs62; mov.b32 {%rs63, %rs64}, %r297; cvt.f32.f16 %r976, %rs63; cvt.f32.f16 %r977, %rs64; fma.rn.f32 %r1181, %r1181, %r901, %r977; fma.rn.f32 %r1180, %r1180, %r901, %r976; fma.rn.f32 %r1179, %r1179, %r900, %r975; fma.rn.f32 %r1178, %r1178, %r900, %r974; fma.rn.f32 %r1185, %r1185, %r901, %r973; fma.rn.f32 %r1184, %r1184, %r901, %r972; fma.rn.f32 %r1183, %r1183, %r900, %r971; fma.rn.f32 %r1182, %r1182, %r900, %r970; fma.rn.f32 %r1189, %r1189, %r901, %r969; fma.rn.f32 %r1188, %r1188, %r901, %r968; fma.rn.f32 %r1187, %r1187, %r900, %r967; fma.rn.f32 %r1186, %r1186, %r900, %r966; fma.rn.f32 %r1193, %r1193, %r901, %r965; fma.rn.f32 %r1192, %r1192, %r901, %r964; fma.rn.f32 %r1191, %r1191, %r900, %r963; fma.rn.f32 %r1190, %r1190, %r900, %r962; fma.rn.f32 %r1197, %r1197, %r901, %r961; fma.rn.f32 %r1196, %r1196, %r901, %r960; fma.rn.f32 %r1195, %r1195, %r900, %r959; fma.rn.f32 %r1194, %r1194, %r900, %r958; fma.rn.f32 %r1201, %r1201, %r901, %r957; fma.rn.f32 %r1200, %r1200, %r901, %r956; fma.rn.f32 %r1199, %r1199, %r900, %r955; fma.rn.f32 %r1198, %r1198, %r900, %r954; fma.rn.f32 %r1205, %r1205, %r901, %r953; fma.rn.f32 %r1204, %r1204, %r901, %r952; fma.rn.f32 %r1203, %r1203, %r900, %r951; fma.rn.f32 %r1202, %r1202, %r900, %r950; fma.rn.f32 %r1209, %r1209, %r901, %r949; fma.rn.f32 %r1208, %r1208, %r901, %r948; fma.rn.f32 %r1207, %r1207, %r900, %r947; fma.rn.f32 %r1206, %r1206, %r900, %r946; fma.rn.f32 %r1213, %r1213, %r903, %r945; fma.rn.f32 %r1212, %r1212, %r903, %r944; fma.rn.f32 %r1211, %r1211, %r902, %r943; fma.rn.f32 %r1210, %r1210, %r902, %r942; fma.rn.f32 %r1217, %r1217, %r903, %r941; fma.rn.f32 %r1216, %r1216, %r903, %r940; fma.rn.f32 %r1215, %r1215, %r902, %r939; fma.rn.f32 %r1214, %r1214, %r902, %r938; fma.rn.f32 %r1221, %r1221, %r903, %r937; fma.rn.f32 %r1220, %r1220, %r903, %r936; fma.rn.f32 %r1219, %r1219, %r902, %r935; fma.rn.f32 %r1218, %r1218, %r902, %r934; fma.rn.f32 %r1225, %r1225, %r903, %r933; fma.rn.f32 %r1224, %r1224, %r903, %r932; fma.rn.f32 %r1223, %r1223, %r902, %r931; fma.rn.f32 %r1222, %r1222, %r902, %r930; fma.rn.f32 %r1229, %r1229, %r903, %r929; fma.rn.f32 %r1228, %r1228, %r903, %r928; fma.rn.f32 %r1227, %r1227, %r902, %r927; fma.rn.f32 %r1226, %r1226, %r902, %r926; fma.rn.f32 %r1233, %r1233, %r903, %r925; fma.rn.f32 %r1232, %r1232, %r903, %r924; fma.rn.f32 %r1231, %r1231, %r902, %r923; fma.rn.f32 %r1230, %r1230, %r902, %r922; fma.rn.f32 %r1237, %r1237, %r903, %r921; fma.rn.f32 %r1236, %r1236, %r903, %r920; fma.rn.f32 %r1235, %r1235, %r902, %r919; fma.rn.f32 %r1234, %r1234, %r902, %r918; fma.rn.f32 %r1241, %r1241, %r903, %r917; fma.rn.f32 %r1240, %r1240, %r903, %r916; fma.rn.f32 %r1239, %r1239, %r902, %r915; fma.rn.f32 %r1238, %r1238, %r902, %r914; .loc 1 34 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:34:9 add.s64 %rd94, %rd126, %rd7; .loc 1 35 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:35:9 add.s64 %rd95, %rd126, %rd8; .loc 1 36 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:36:9 add.s64 %rd96, %rd127, %rd22; add.s64 %rd97, %rd127, %rd21; add.s64 %rd98, %rd127, %rd20; .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 add.s64 %rd99, %rd127, %rd18; add.s32 %r978, %r1171, 1; setp.gt.s32 %p22, %r978, 1; selp.b32 %r1171, 0, %r978, %p22; .loc 1 15 18, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:15:18 setp.lt.s32 %p23, %r3, %r1169; setp.lt.s32 %p24, %r4, %r1169; .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 shl.b32 %r979, %r1171, 12; add.s32 %r980, %r126, %r979; bar.sync 0; add.s32 %r981, %r13, %r979; add.s32 %r400, %r981, 24576; selp.b32 %r982, 16, 0, %p23; selp.b32 %r401, %r982, 0, %p20; // begin inline asm .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.cg.shared.global [ %r400 + 0 ], [ %rd94 + 0 ], 0x10, %r401; // end inline asm add.s32 %r402, %r981, 26624; selp.b32 %r983, 16, 0, %p24; selp.b32 %r403, %r983, 0, %p20; // begin inline asm .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 cp.async.cg.shared.global [ %r402 + 0 ], [ %rd95 + 0 ], 0x10, %r403; // end inline asm cp.async.commit_group; .loc 1 17 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:17:19 shl.b32 %r984, %r1171, 2; add.s32 %r985, %r126, %r984; add.s32 %r404, %r985, 32768; selp.b32 %r405, 4, 0, %p20; // begin inline asm @%p5 cp.async.ca.shared.global [ %r404 + 0 ], [ %rd125 + 0 ], 0x4, %r405; // end inline asm cp.async.commit_group; .loc 1 29 36, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:36 setp.lt.s32 %p25, %r6, %r1169; setp.lt.s32 %p26, %r7, %r1169; setp.lt.s32 %p27, %r8, %r1169; setp.lt.s32 %p28, %r9, %r1169; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 add.s32 %r986, %r980, %r979; add.s32 %r406, %r986, %r15; selp.b32 %r987, 16, 0, %p25; selp.b32 %r407, %r987, 0, %p20; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r406 + 0 ], [ %rd96 + 0 ], 0x10, %r407; // end inline asm add.s32 %r408, %r406, 2048; selp.b32 %r988, 16, 0, %p26; selp.b32 %r409, %r988, 0, %p20; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r408 + 0 ], [ %rd97 + 0 ], 0x10, %r409; // end inline asm add.s32 %r410, %r406, 4096; selp.b32 %r989, 16, 0, %p27; selp.b32 %r411, %r989, 0, %p20; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r410 + 0 ], [ %rd98 + 0 ], 0x10, %r411; // end inline asm add.s32 %r412, %r406, 6144; selp.b32 %r990, 16, 0, %p28; selp.b32 %r413, %r990, 0, %p20; // begin inline asm .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 cp.async.cg.shared.global [ %r412 + 0 ], [ %rd99 + 0 ], 0x10, %r413; // end inline asm cp.async.commit_group; .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 add.s32 %r1172, %r1172, 64; add.s64 %rd127, %rd127, %rd19; add.s64 %rd126, %rd126, %rd16; add.s64 %rd125, %rd125, 4; add.s32 %r1169, %r1169, -64; setp.lt.s32 %p29, %r1172, %r38; mov.b32 %r1173, %r30; mov.b32 %r1174, %r31; mov.b32 %r1175, %r32; mov.b32 %r1176, %r33; @%p29 bra $L__BB0_3; bra.uni $L__BB0_4; $L__tmp30: $L__BB0_1: // %.._crit_edge_crit_edge .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 and.b32 %r1177, %r2, 16; mov.b32 %r1242, 0f3F800000; mov.b32 %r1178, 0f00000000; mov.b32 %r1179, %r1178; mov.b32 %r1180, %r1178; mov.b32 %r1181, %r1178; mov.b32 %r1182, %r1178; mov.b32 %r1183, %r1178; mov.b32 %r1184, %r1178; mov.b32 %r1185, %r1178; mov.b32 %r1186, %r1178; mov.b32 %r1187, %r1178; mov.b32 %r1188, %r1178; mov.b32 %r1189, %r1178; mov.b32 %r1190, %r1178; mov.b32 %r1191, %r1178; mov.b32 %r1192, %r1178; mov.b32 %r1193, %r1178; mov.b32 %r1194, %r1178; mov.b32 %r1195, %r1178; mov.b32 %r1196, %r1178; mov.b32 %r1197, %r1178; mov.b32 %r1198, %r1178; mov.b32 %r1199, %r1178; mov.b32 %r1200, %r1178; mov.b32 %r1201, %r1178; mov.b32 %r1202, %r1178; mov.b32 %r1203, %r1178; mov.b32 %r1204, %r1178; mov.b32 %r1205, %r1178; mov.b32 %r1206, %r1178; mov.b32 %r1207, %r1178; mov.b32 %r1208, %r1178; mov.b32 %r1209, %r1178; mov.b32 %r1210, %r1178; mov.b32 %r1211, %r1178; mov.b32 %r1212, %r1178; mov.b32 %r1213, %r1178; mov.b32 %r1214, %r1178; mov.b32 %r1215, %r1178; mov.b32 %r1216, %r1178; mov.b32 %r1217, %r1178; mov.b32 %r1218, %r1178; mov.b32 %r1219, %r1178; mov.b32 %r1220, %r1178; mov.b32 %r1221, %r1178; mov.b32 %r1222, %r1178; mov.b32 %r1223, %r1178; mov.b32 %r1224, %r1178; mov.b32 %r1225, %r1178; mov.b32 %r1226, %r1178; mov.b32 %r1227, %r1178; mov.b32 %r1228, %r1178; mov.b32 %r1229, %r1178; mov.b32 %r1230, %r1178; mov.b32 %r1231, %r1178; mov.b32 %r1232, %r1178; mov.b32 %r1233, %r1178; mov.b32 %r1234, %r1178; mov.b32 %r1235, %r1178; mov.b32 %r1236, %r1178; mov.b32 %r1237, %r1178; mov.b32 %r1238, %r1178; mov.b32 %r1239, %r1178; mov.b32 %r1240, %r1178; mov.b32 %r1241, %r1178; mov.b32 %r1243, %r1242; mov.b32 %r1244, %r1242; mov.b32 %r1245, %r1242; $L__BB0_4: // %._crit_edge .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1059, %r5, %r1; or.b32 %r1060, %r1059, 112; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p37, %r1060, %r37; .loc 1 59 34 // attn_qk_int8_per_block.py:59:34 or.b32 %r1061, %r6, %r1; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1062, %r1061, 96; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p36, %r1062, %r37; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1063, %r1061, 80; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p35, %r1063, %r37; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1064, %r1061, 64; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p34, %r1064, %r37; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1065, %r9, %r1; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p33, %r1065, %r37; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1066, %r8, %r1; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p32, %r1066, %r37; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r1067, %r7, %r1; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p31, %r1067, %r37; setp.lt.s32 %p30, %r1061, %r37; .loc 1 67 26 // attn_qk_int8_per_block.py:67:26 cvt.s64.s32 %rd108, %r34; mul.lo.s64 %rd109, %rd108, %rd1; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 shl.b64 %rd110, %rd109, 1; add.s64 %rd111, %rd26, %rd110; .loc 1 67 46 // attn_qk_int8_per_block.py:67:46 cvt.s64.s32 %rd112, %r35; mul.lo.s64 %rd113, %rd112, %rd2; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 shl.b64 %rd114, %rd113, 1; add.s64 %rd115, %rd111, %rd114; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1068, %r36, %r1060; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd116, %r1068, 2, %rd115; shl.b64 %rd117, %rd15, 1; add.s64 %rd107, %rd116, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1069, %r36, %r1062; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd118, %r1069, 2, %rd115; add.s64 %rd106, %rd118, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1070, %r36, %r1063; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd119, %r1070, 2, %rd115; add.s64 %rd105, %rd119, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1071, %r36, %r1064; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd120, %r1071, 2, %rd115; add.s64 %rd104, %rd120, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1072, %r36, %r1065; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd121, %r1072, 2, %rd115; add.s64 %rd103, %rd121, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1073, %r36, %r1066; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd122, %r1073, 2, %rd115; add.s64 %rd102, %rd122, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1074, %r36, %r1067; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd123, %r1074, 2, %rd115; add.s64 %rd101, %rd123, %rd117; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r1075, %r36, %r1061; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd124, %r1075, 2, %rd115; add.s64 %rd100, %rd124, %rd117; $L__tmp31: .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 cp.async.wait_group 0; bar.sync 0; $L__tmp32: .loc 1 80 11 // attn_qk_int8_per_block.py:80:11 div.full.f32 %r1076, %r1178, %r1242; div.full.f32 %r1077, %r1179, %r1242; div.full.f32 %r1078, %r1180, %r1243; div.full.f32 %r1079, %r1181, %r1243; div.full.f32 %r1080, %r1182, %r1242; div.full.f32 %r1081, %r1183, %r1242; div.full.f32 %r1082, %r1184, %r1243; div.full.f32 %r1083, %r1185, %r1243; div.full.f32 %r1084, %r1186, %r1242; div.full.f32 %r1085, %r1187, %r1242; div.full.f32 %r1086, %r1188, %r1243; div.full.f32 %r1087, %r1189, %r1243; div.full.f32 %r1088, %r1190, %r1242; div.full.f32 %r1089, %r1191, %r1242; div.full.f32 %r1090, %r1192, %r1243; div.full.f32 %r1091, %r1193, %r1243; div.full.f32 %r1092, %r1194, %r1242; div.full.f32 %r1093, %r1195, %r1242; div.full.f32 %r1094, %r1196, %r1243; div.full.f32 %r1095, %r1197, %r1243; div.full.f32 %r1096, %r1198, %r1242; div.full.f32 %r1097, %r1199, %r1242; div.full.f32 %r1098, %r1200, %r1243; div.full.f32 %r1099, %r1201, %r1243; div.full.f32 %r1100, %r1202, %r1242; div.full.f32 %r1101, %r1203, %r1242; div.full.f32 %r1102, %r1204, %r1243; div.full.f32 %r1103, %r1205, %r1243; div.full.f32 %r1104, %r1206, %r1242; div.full.f32 %r1105, %r1207, %r1242; div.full.f32 %r1106, %r1208, %r1243; div.full.f32 %r1107, %r1209, %r1243; div.full.f32 %r1108, %r1210, %r1244; div.full.f32 %r1109, %r1211, %r1244; div.full.f32 %r1110, %r1212, %r1245; div.full.f32 %r1111, %r1213, %r1245; div.full.f32 %r1112, %r1214, %r1244; div.full.f32 %r1113, %r1215, %r1244; div.full.f32 %r1114, %r1216, %r1245; div.full.f32 %r1115, %r1217, %r1245; div.full.f32 %r1116, %r1218, %r1244; div.full.f32 %r1117, %r1219, %r1244; div.full.f32 %r1118, %r1220, %r1245; div.full.f32 %r1119, %r1221, %r1245; div.full.f32 %r1120, %r1222, %r1244; div.full.f32 %r1121, %r1223, %r1244; div.full.f32 %r1122, %r1224, %r1245; div.full.f32 %r1123, %r1225, %r1245; div.full.f32 %r1124, %r1226, %r1244; div.full.f32 %r1125, %r1227, %r1244; div.full.f32 %r1126, %r1228, %r1245; div.full.f32 %r1127, %r1229, %r1245; div.full.f32 %r1128, %r1230, %r1244; div.full.f32 %r1129, %r1231, %r1244; div.full.f32 %r1130, %r1232, %r1245; div.full.f32 %r1131, %r1233, %r1245; div.full.f32 %r1132, %r1234, %r1244; div.full.f32 %r1133, %r1235, %r1244; div.full.f32 %r1134, %r1236, %r1245; div.full.f32 %r1135, %r1237, %r1245; div.full.f32 %r1136, %r1238, %r1244; div.full.f32 %r1137, %r1239, %r1244; div.full.f32 %r1138, %r1240, %r1245; div.full.f32 %r1139, %r1241, %r1245; .loc 1 81 27 // attn_qk_int8_per_block.py:81:27 cvt.rn.bf16x2.f32 %r992, %r1077, %r1076; cvt.rn.bf16x2.f32 %r997, %r1079, %r1078; cvt.rn.bf16x2.f32 %r993, %r1081, %r1080; cvt.rn.bf16x2.f32 %r998, %r1083, %r1082; cvt.rn.bf16x2.f32 %r994, %r1085, %r1084; cvt.rn.bf16x2.f32 %r999, %r1087, %r1086; cvt.rn.bf16x2.f32 %r995, %r1089, %r1088; cvt.rn.bf16x2.f32 %r1000, %r1091, %r1090; cvt.rn.bf16x2.f32 %r1002, %r1093, %r1092; cvt.rn.bf16x2.f32 %r1007, %r1095, %r1094; cvt.rn.bf16x2.f32 %r1003, %r1097, %r1096; cvt.rn.bf16x2.f32 %r1008, %r1099, %r1098; cvt.rn.bf16x2.f32 %r1004, %r1101, %r1100; cvt.rn.bf16x2.f32 %r1009, %r1103, %r1102; cvt.rn.bf16x2.f32 %r1005, %r1105, %r1104; cvt.rn.bf16x2.f32 %r1010, %r1107, %r1106; cvt.rn.bf16x2.f32 %r1011, %r1109, %r1108; cvt.rn.bf16x2.f32 %r1015, %r1111, %r1110; cvt.rn.bf16x2.f32 %r1012, %r1113, %r1112; cvt.rn.bf16x2.f32 %r1016, %r1115, %r1114; cvt.rn.bf16x2.f32 %r1013, %r1117, %r1116; cvt.rn.bf16x2.f32 %r1017, %r1119, %r1118; cvt.rn.bf16x2.f32 %r1014, %r1121, %r1120; cvt.rn.bf16x2.f32 %r1018, %r1123, %r1122; cvt.rn.bf16x2.f32 %r1019, %r1125, %r1124; cvt.rn.bf16x2.f32 %r1023, %r1127, %r1126; cvt.rn.bf16x2.f32 %r1020, %r1129, %r1128; cvt.rn.bf16x2.f32 %r1024, %r1131, %r1130; cvt.rn.bf16x2.f32 %r1021, %r1133, %r1132; cvt.rn.bf16x2.f32 %r1025, %r1135, %r1134; cvt.rn.bf16x2.f32 %r1022, %r1137, %r1136; cvt.rn.bf16x2.f32 %r1026, %r1139, %r1138; .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 and.b32 %r1140, %r10, 1536; and.b32 %r1141, %r2, 1; shl.b32 %r1142, %r1141, 5; and.b32 %r1143, %r2, 2; bfe.s32 %r1144, %r2, 1, 1; and.b32 %r1145, %r1144, 4160; bfe.s32 %r1146, %r2, 2, 1; and.b32 %r1147, %r2, 4; shl.b32 %r1148, %r1147, 2; bfe.s32 %r1149, %r2, 3, 1; and.b32 %r1150, %r1149, 2080; shl.b32 %r1151, %r1177, 3; or.b32 %r1152, %r1145, %r1148; or.b32 %r1153, %r1152, %r1151; or.b32 %r1154, %r1140, %r1142; xor.b32 %r1155, %r1154, %r1150; or.b32 %r1156, %r1153, %r1155; add.s32 %r991, %r126, %r1156; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r991 + 0 ], { %r992, %r993, %r994, %r995 }; // end inline asm add.s32 %r996, %r991, 256; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r996 + 0 ], { %r997, %r998, %r999, %r1000 }; // end inline asm xor.b32 %r1157, %r1156, 64; add.s32 %r1001, %r126, %r1157; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r1001 + 0 ], { %r1002, %r1003, %r1004, %r1005 }; // end inline asm add.s32 %r1006, %r1001, 256; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r1006 + 0 ], { %r1007, %r1008, %r1009, %r1010 }; // end inline asm bar.sync 0; shl.b32 %r1158, %r2, 2; and.b32 %r1159, %r1158, 416; shl.b32 %r1160, %r1141, 6; shl.b32 %r1161, %r1143, 3; and.b32 %r1162, %r1146, 2080; setp.eq.b32 %p38, %r1177, 0; selp.b32 %r1163, 0, 4160, %p38; or.b32 %r1164, %r1159, %r1160; or.b32 %r1165, %r1162, %r1163; xor.b32 %r1166, %r1165, %r1164; add.s32 %r1167, %r126, %r1161; add.s32 %r1168, %r1167, %r1166; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1027, %r1028, %r1029, %r1030}, [%r1168]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1031, %r1032, %r1033, %r1034}, [%r1168+512]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1035, %r1036, %r1037, %r1038}, [%r1168+1024]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1039, %r1040, %r1041, %r1042}, [%r1168+1536]; bar.sync 0; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r991 + 0 ], { %r1011, %r1012, %r1013, %r1014 }; // end inline asm // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r996 + 0 ], { %r1015, %r1016, %r1017, %r1018 }; // end inline asm // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r1001 + 0 ], { %r1019, %r1020, %r1021, %r1022 }; // end inline asm // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r1006 + 0 ], { %r1023, %r1024, %r1025, %r1026 }; // end inline asm bar.sync 0; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1043, %r1044, %r1045, %r1046}, [%r1168]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1047, %r1048, %r1049, %r1050}, [%r1168+512]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1051, %r1052, %r1053, %r1054}, [%r1168+1024]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r1055, %r1056, %r1057, %r1058}, [%r1168+1536]; // begin inline asm @%p30 st.global.v4.b32 [ %rd100 + 0 ], { %r1027, %r1028, %r1029, %r1030 }; // end inline asm // begin inline asm @%p31 st.global.v4.b32 [ %rd101 + 0 ], { %r1031, %r1032, %r1033, %r1034 }; // end inline asm // begin inline asm @%p32 st.global.v4.b32 [ %rd102 + 0 ], { %r1035, %r1036, %r1037, %r1038 }; // end inline asm // begin inline asm @%p33 st.global.v4.b32 [ %rd103 + 0 ], { %r1039, %r1040, %r1041, %r1042 }; // end inline asm // begin inline asm @%p34 st.global.v4.b32 [ %rd104 + 0 ], { %r1043, %r1044, %r1045, %r1046 }; // end inline asm // begin inline asm @%p35 st.global.v4.b32 [ %rd105 + 0 ], { %r1047, %r1048, %r1049, %r1050 }; // end inline asm // begin inline asm @%p36 st.global.v4.b32 [ %rd106 + 0 ], { %r1051, %r1052, %r1053, %r1054 }; // end inline asm // begin inline asm @%p37 st.global.v4.b32 [ %rd107 + 0 ], { %r1055, %r1056, %r1057, %r1058 }; // end inline asm .loc 1 40 1 // attn_qk_int8_per_block.py:40:1 ret; $L__tmp33: $L__func_end0: // -- End function } .file 1 "D:\\Comfy\\App\\ComfyUI_windows_portable\\python_embeded\\Lib\\site-packages\\sageattention\\attn_qk_int8_per_block.py" .file 2 "D:\\Comfy\\App\\ComfyUI_windows_portable\\python_embeded\\Lib\\site-packages\\triton\\language\\standard.py" .section .debug_abbrev { .b8 1 // Abbreviation Code .b8 17 // DW_TAG_compile_unit .b8 1 // DW_CHILDREN_yes .b8 37 // DW_AT_producer .b8 8 // DW_FORM_string .b8 19 // DW_AT_language .b8 5 // DW_FORM_data2 .b8 3 // DW_AT_name .b8 8 // DW_FORM_string .b8 16 // DW_AT_stmt_list .b8 6 // DW_FORM_data4 .b8 27 // DW_AT_comp_dir .b8 8 // DW_FORM_string .b8 0 // EOM(1) .b8 0 // EOM(2) .b8 2 // Abbreviation Code .b8 46 // DW_TAG_subprogram .b8 0 // DW_CHILDREN_no .b8 3 // DW_AT_name .b8 8 // DW_FORM_string .b8 32 // DW_AT_inline .b8 11 // DW_FORM_data1 .b8 0 // EOM(1) .b8 0 // EOM(2) .b8 3 // Abbreviation Code .b8 46 // DW_TAG_subprogram .b8 1 // DW_CHILDREN_yes .b8 17 // DW_AT_low_pc .b8 1 // DW_FORM_addr .b8 18 // DW_AT_high_pc .b8 1 // DW_FORM_addr .b8 49 // DW_AT_abstract_origin .b8 19 // DW_FORM_ref4 .b8 0 // EOM(1) .b8 0 // EOM(2) .b8 4 // Abbreviation Code .b8 29 // DW_TAG_inlined_subroutine .b8 0 // DW_CHILDREN_no .b8 49 // DW_AT_abstract_origin .b8 19 // DW_FORM_ref4 .b8 17 // DW_AT_low_pc .b8 1 // DW_FORM_addr .b8 18 // DW_AT_high_pc .b8 1 // DW_FORM_addr .b8 88 // DW_AT_call_file .b8 11 // DW_FORM_data1 .b8 89 // DW_AT_call_line .b8 11 // DW_FORM_data1 .b8 87 // DW_AT_call_column .b8 11 // DW_FORM_data1 .b8 0 // EOM(1) .b8 0 // EOM(2) .b8 5 // Abbreviation Code .b8 29 // DW_TAG_inlined_subroutine .b8 1 // DW_CHILDREN_yes .b8 49 // DW_AT_abstract_origin .b8 19 // DW_FORM_ref4 .b8 17 // DW_AT_low_pc .b8 1 // DW_FORM_addr .b8 18 // DW_AT_high_pc .b8 1 // DW_FORM_addr .b8 88 // DW_AT_call_file .b8 11 // DW_FORM_data1 .b8 89 // DW_AT_call_line .b8 11 // DW_FORM_data1 .b8 87 // DW_AT_call_column .b8 11 // DW_FORM_data1 .b8 0 // EOM(1) .b8 0 // EOM(2) .b8 6 // Abbreviation Code .b8 29 // DW_TAG_inlined_subroutine .b8 0 // DW_CHILDREN_no .b8 49 // DW_AT_abstract_origin .b8 19 // DW_FORM_ref4 .b8 17 // DW_AT_low_pc .b8 1 // DW_FORM_addr .b8 18 // DW_AT_high_pc .b8 1 // DW_FORM_addr .b8 88 // DW_AT_call_file .b8 11 // DW_FORM_data1 .b8 89 // DW_AT_call_line .b8 5 // DW_FORM_data2 .b8 87 // DW_AT_call_column .b8 11 // DW_FORM_data1 .b8 0 // EOM(1) .b8 0 // EOM(2) .b8 0 // EOM(3) } .section .debug_info { .b32 315 // Length of Unit .b8 2 // DWARF version number .b8 0 .b32 .debug_abbrev // Offset Into Abbrev. Section .b8 8 // Address Size (in bytes) .b8 1 // Abbrev [1] 0xb:0x134 DW_TAG_compile_unit .b8 116 // DW_AT_producer .b8 114 .b8 105 .b8 116 .b8 111 .b8 110 .b8 0 .b8 2 // DW_AT_language .b8 0 .b8 97 // DW_AT_name .b8 116 .b8 116 .b8 110 .b8 95 .b8 113 .b8 107 .b8 95 .b8 105 .b8 110 .b8 116 .b8 56 .b8 95 .b8 112 .b8 101 .b8 114 .b8 95 .b8 98 .b8 108 .b8 111 .b8 99 .b8 107 .b8 46 .b8 112 .b8 121 .b8 0 .b32 .debug_line // DW_AT_stmt_list .b8 68 // DW_AT_comp_dir .b8 58 .b8 92 .b8 67 .b8 111 .b8 109 .b8 102 .b8 121 .b8 92 .b8 65 .b8 112 .b8 112 .b8 92 .b8 67 .b8 111 .b8 109 .b8 102 .b8 121 .b8 85 .b8 73 .b8 95 .b8 119 .b8 105 .b8 110 .b8 100 .b8 111 .b8 119 .b8 115 .b8 95 .b8 112 .b8 111 .b8 114 .b8 116 .b8 97 .b8 98 .b8 108 .b8 101 .b8 92 .b8 112 .b8 121 .b8 116 .b8 104 .b8 111 .b8 110 .b8 95 .b8 101 .b8 109 .b8 98 .b8 101 .b8 100 .b8 101 .b8 100 .b8 92 .b8 76 .b8 105 .b8 98 .b8 92 .b8 115 .b8 105 .b8 116 .b8 101 .b8 45 .b8 112 .b8 97 .b8 99 .b8 107 .b8 97 .b8 103 .b8 101 .b8 115 .b8 92 .b8 115 .b8 97 .b8 103 .b8 101 .b8 97 .b8 116 .b8 116 .b8 101 .b8 110 .b8 116 .b8 105 .b8 111 .b8 110 .b8 0 .b8 2 // Abbrev [2] 0x88:0xc DW_TAG_subprogram .b8 95 // DW_AT_name .b8 97 .b8 116 .b8 116 .b8 110 .b8 95 .b8 102 .b8 119 .b8 100 .b8 0 .b8 1 // DW_AT_inline .b8 3 // Abbrev [3] 0x94:0xaa DW_TAG_subprogram .b64 $L__func_begin0 // DW_AT_low_pc .b64 $L__func_end0 // DW_AT_high_pc .b32 136 // DW_AT_abstract_origin .b8 4 // Abbrev [4] 0xa9:0x18 DW_TAG_inlined_subroutine .b32 136 // DW_AT_abstract_origin .b64 $L__tmp1 // DW_AT_low_pc .b64 $L__tmp2 // DW_AT_high_pc .b8 1 // DW_AT_call_file .b8 56 // DW_AT_call_line .b8 44 // DW_AT_call_column .b8 5 // Abbrev [5] 0xc1:0x7c DW_TAG_inlined_subroutine .b32 136 // DW_AT_abstract_origin .b64 $L__tmp3 // DW_AT_low_pc .b64 $L__tmp32 // DW_AT_high_pc .b8 1 // DW_AT_call_file .b8 75 // DW_AT_call_line .b8 16 // DW_AT_call_column .b8 5 // Abbrev [5] 0xd9:0x31 DW_TAG_inlined_subroutine .b32 136 // DW_AT_abstract_origin .b64 $L__tmp6 // DW_AT_low_pc .b64 $L__tmp17 // DW_AT_high_pc .b8 1 // DW_AT_call_file .b8 19 // DW_AT_call_line .b8 32 // DW_AT_call_column .b8 4 // Abbrev [4] 0xf1:0x18 DW_TAG_inlined_subroutine .b32 136 // DW_AT_abstract_origin .b64 $L__tmp6 // DW_AT_low_pc .b64 $L__tmp17 // DW_AT_high_pc .b8 2 // DW_AT_call_file .b8 191 // DW_AT_call_line .b8 16 // DW_AT_call_column .b8 0 // End Of Children Mark .b8 5 // Abbrev [5] 0x10a:0x32 DW_TAG_inlined_subroutine .b32 136 // DW_AT_abstract_origin .b64 $L__tmp18 // DW_AT_low_pc .b64 $L__tmp29 // DW_AT_high_pc .b8 1 // DW_AT_call_file .b8 22 // DW_AT_call_line .b8 16 // DW_AT_call_column .b8 6 // Abbrev [6] 0x122:0x19 DW_TAG_inlined_subroutine .b32 136 // DW_AT_abstract_origin .b64 $L__tmp18 // DW_AT_low_pc .b64 $L__tmp29 // DW_AT_high_pc .b8 2 // DW_AT_call_file .b8 37 // DW_AT_call_line .b8 1 .b8 12 // DW_AT_call_column .b8 0 // End Of Children Mark .b8 0 // End Of Children Mark .b8 0 // End Of Children Mark .b8 0 // End Of Children Mark } .section .debug_str { $L__info_string0: .b8 95 // string offset=0 ; _attn_fwd .b8 97 .b8 116 .b8 116 .b8 110 .b8 95 .b8 102 .b8 119 .b8 100 .b8 0 } .section .debug_macinfo { }