// // 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 256 { .reg .pred %p<46>; .reg .b16 %rs<65>; .reg .b32 %r<1067>; .reg .b64 %rd<137>; .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 %r37, [_attn_fwd_param_19]; ld.param.b32 %r36, [_attn_fwd_param_18]; ld.param.b32 %r35, [_attn_fwd_param_17]; ld.param.b32 %r34, [_attn_fwd_param_16]; ld.param.b32 %r33, [_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 %rd53, [_attn_fwd_param_0]; $L__tmp0: .loc 1 51 15 // attn_qk_int8_per_block.py:51:15 mov.u32 %r101, %ctaid.x; .loc 1 53 13 // attn_qk_int8_per_block.py:53:13 mov.u32 %r102, %ctaid.z; cvt.u64.u32 %rd1, %r102; ld.param.b64 %rd54, [_attn_fwd_param_3]; .loc 1 54 13 // attn_qk_int8_per_block.py:54:13 mov.u32 %r103, %ctaid.y; cvt.u64.u32 %rd2, %r103; ld.param.s32 %rd55, [_attn_fwd_param_6]; .loc 1 56 23 // attn_qk_int8_per_block.py:56:23 mad.wide.u32 %rd56, %r102, 32, %rd2; ld.param.s32 %rd57, [_attn_fwd_param_7]; ld.param.b32 %r104, [_attn_fwd_param_8]; ld.param.s32 %rd58, [_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 %r105, %r36, 127; add.s32 %r106, %r37, 63; ld.param.s32 %rd59, [_attn_fwd_param_10]; .loc 2 43 12, function_name $L__info_string0, inlined_at 1 56 44 // standard.py:43:12 shr.s32 %r107, %r106, 31; shr.u32 %r108, %r107, 26; add.s32 %r109, %r106, %r108; shr.s32 %r110, %r109, 6; shr.s32 %r111, %r105, 31; shr.u32 %r112, %r111, 25; add.s32 %r113, %r105, %r112; shr.s32 %r114, %r113, 7; ld.param.b32 %r115, [_attn_fwd_param_11]; ld.param.s32 %rd60, [_attn_fwd_param_12]; $L__tmp2: .loc 1 56 22 // attn_qk_int8_per_block.py:56:22 cvt.s64.s32 %rd61, %r114; ld.param.s32 %rd62, [_attn_fwd_param_13]; mul.lo.s64 %rd63, %rd56, %rd61; ld.param.b32 %r116, [_attn_fwd_param_14]; .loc 1 57 22 // attn_qk_int8_per_block.py:57:22 cvt.s64.s32 %rd64, %r110; mul.lo.s64 %rd3, %rd56, %rd64; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 shl.b32 %r1, %r101, 7; .loc 1 59 34 // attn_qk_int8_per_block.py:59:34 mov.u32 %r2, %tid.x; shr.u32 %r117, %r2, 3; bfe.u32 %r3, %r2, 3, 5; or.b32 %r4, %r3, 32; shr.u32 %r5, %r2, 4; bfe.u32 %r6, %r2, 4, 4; or.b32 %r7, %r6, 16; or.b32 %r8, %r6, 32; or.b32 %r9, %r5, 48; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r118, %r3, %r1; or.b32 %r119, %r4, %r1; or.b32 %r120, %r118, 64; or.b32 %r121, %r1, %r117; or.b32 %r122, %r121, 96; .loc 1 62 14 // attn_qk_int8_per_block.py:62:14 mad.lo.s64 %rd65, %rd55, %rd1, %rd53; mad.lo.s64 %rd66, %rd57, %rd2, %rd65; .loc 1 62 60 // attn_qk_int8_per_block.py:62:60 mul.lo.s32 %r123, %r104, %r118; mul.lo.s32 %r124, %r104, %r119; mul.lo.s32 %r125, %r104, %r120; mul.lo.s32 %r126, %r104, %r122; .loc 1 62 14 // attn_qk_int8_per_block.py:62:14 cvt.s64.s32 %rd67, %r123; add.s64 %rd68, %rd66, %rd67; cvt.s64.s32 %rd69, %r124; add.s64 %rd70, %rd66, %rd69; cvt.s64.s32 %rd71, %r125; add.s64 %rd72, %rd66, %rd71; cvt.s64.s32 %rd73, %r126; add.s64 %rd74, %rd66, %rd73; .loc 1 62 90 // attn_qk_int8_per_block.py:62:90 and.b32 %r10, %r2, 7; shl.b32 %r127, %r10, 4; and.b32 %r11, %r2, 15; shl.b32 %r128, %r11, 3; .loc 1 62 14 // attn_qk_int8_per_block.py:62:14 cvt.u64.u32 %rd4, %r127; add.s64 %rd27, %rd68, %rd4; add.s64 %rd28, %rd70, %rd4; add.s64 %rd29, %rd72, %rd4; add.s64 %rd30, %rd74, %rd4; .loc 1 63 19 // attn_qk_int8_per_block.py:63:19 shl.b64 %rd75, %rd63, 2; add.s64 %rd76, %rd54, %rd75; mad.wide.u32 %rd31, %r101, 4, %rd76; .loc 1 64 19 // attn_qk_int8_per_block.py:64:19 mul.lo.s64 %rd5, %rd58, %rd1; .loc 1 64 39 // attn_qk_int8_per_block.py:64:39 mul.lo.s64 %rd6, %rd59, %rd2; .loc 1 64 14 // attn_qk_int8_per_block.py:64:14 add.s64 %rd77, %rd23, %rd5; add.s64 %rd78, %rd77, %rd6; .loc 1 64 79 // attn_qk_int8_per_block.py:64:79 mul.lo.s32 %r129, %r115, %r3; shl.b32 %r130, %r115, 5; add.s32 %r131, %r129, %r130; .loc 1 64 14 // attn_qk_int8_per_block.py:64:14 cvt.s64.s32 %rd7, %r129; add.s64 %rd79, %rd78, %rd7; cvt.s64.s32 %rd8, %r131; add.s64 %rd80, %rd78, %rd8; add.s64 %rd32, %rd79, %rd4; add.s64 %rd33, %rd80, %rd4; .loc 1 65 19 // attn_qk_int8_per_block.py:65:19 shl.b64 %rd81, %rd3, 2; add.s64 %rd34, %rd25, %rd81; .loc 1 66 19 // attn_qk_int8_per_block.py:66:19 mul.lo.s64 %rd9, %rd60, %rd1; .loc 1 66 39 // attn_qk_int8_per_block.py:66:39 mul.lo.s64 %rd10, %rd62, %rd2; .loc 1 66 14 // attn_qk_int8_per_block.py:66:14 shl.b64 %rd82, %rd9, 1; add.s64 %rd83, %rd24, %rd82; shl.b64 %rd84, %rd10, 1; add.s64 %rd85, %rd83, %rd84; .loc 1 66 79 // attn_qk_int8_per_block.py:66:79 mul.lo.s32 %r132, %r116, %r6; shl.b32 %r133, %r116, 4; add.s32 %r134, %r132, %r133; shl.b32 %r135, %r116, 5; add.s32 %r136, %r132, %r135; mul.lo.s32 %r137, %r116, %r9; .loc 1 66 14 // attn_qk_int8_per_block.py:66:14 mad.wide.s32 %rd86, %r132, 2, %rd85; mad.wide.s32 %rd87, %r134, 2, %rd85; mad.wide.s32 %rd88, %r136, 2, %rd85; mad.wide.s32 %rd89, %r137, 2, %rd85; cvt.u64.u32 %rd15, %r128; mul.wide.u32 %rd90, %r128, 2; add.s64 %rd35, %rd86, %rd90; add.s64 %rd36, %rd87, %rd90; add.s64 %rd37, %rd88, %rd90; add.s64 %rd38, %rd89, %rd90; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p1, %r118, %r36; setp.lt.s32 %p2, %r119, %r36; setp.lt.s32 %p3, %r120, %r36; setp.lt.s32 %p4, %r122, %r36; .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 %r38, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r39, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r40, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r41, 0x0; @%p1 ld.global.v4.b32 { %r38, %r39, %r40, %r41 }, [ %rd27 + 0 ]; // end inline asm // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r42, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r43, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r44, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r45, 0x0; @%p2 ld.global.v4.b32 { %r42, %r43, %r44, %r45 }, [ %rd28 + 0 ]; // end inline asm // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r46, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r47, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r48, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r49, 0x0; @%p3 ld.global.v4.b32 { %r46, %r47, %r48, %r49 }, [ %rd29 + 0 ]; // end inline asm // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 mov.u32 %r50, 0x0; .loc 1 74 9 // attn_qk_int8_per_block.py:74:9 mov.u32 %r51, 0x0; .loc 1 75 9 // attn_qk_int8_per_block.py:75:9 mov.u32 %r52, 0x0; .loc 1 76 9 // attn_qk_int8_per_block.py:76:9 mov.u32 %r53, 0x0; @%p4 ld.global.v4.b32 { %r50, %r51, %r52, %r53 }, [ %rd30 + 0 ]; // end inline asm shl.b32 %r138, %r2, 4; and.b32 %r139, %r138, 4080; shl.b32 %r12, %r2, 1; and.b32 %r140, %r12, 112; xor.b32 %r141, %r139, %r140; mov.b32 %r142, global_smem; add.s32 %r13, %r142, %r141; add.s32 %r54, %r13, 73728; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r54 + 0 ], { %r38, %r39, %r40, %r41 }; // end inline asm add.s32 %r55, %r13, 77824; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r55 + 0 ], { %r42, %r43, %r44, %r45 }; // end inline asm add.s32 %r56, %r13, 81920; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r56 + 0 ], { %r46, %r47, %r48, %r49 }; // end inline asm add.s32 %r57, %r13, 86016; // begin inline asm .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 st.shared::cta.v4.b32 [ %r57 + 0 ], { %r50, %r51, %r52, %r53 }; // 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 %r58, 0x0; .loc 1 75 15 // attn_qk_int8_per_block.py:75:15 ld.global.b32 { %r58 }, [ %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 %r143, %r115, 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 %r144, %r116, 6; .loc 1 13 5, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:13:5 setp.lt.s32 %p6, %r37, 1; setp.gt.s32 %p7, %r37, 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 %p8, %r3, %r37; setp.lt.s32 %p9, %r4, %r37; .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 add.s32 %r59, %r13, 49152; selp.b32 %r60, 16, 0, %p8; // 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 [ %r59 + 0 ], [ %rd32 + 0 ], 0x10, %r60; // end inline asm add.s32 %r61, %r13, 53248; selp.b32 %r145, 16, 0, %p9; selp.b32 %r62, %r145, 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 [ %r61 + 0 ], [ %rd33 + 0 ], 0x10, %r62; // 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, 255; setp.eq.b32 %p5, %r14, 0; selp.b32 %r64, 4, 0, %p7; add.s32 %r63, %r142, 90112; // begin inline asm @%p5 cp.async.ca.shared.global [ %r63 + 0 ], [ %rd34 + 0 ], 0x4, %r64; // 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 %p10, %r6, %r37; setp.lt.s32 %p11, %r7, %r37; setp.lt.s32 %p12, %r8, %r37; setp.lt.s32 %p13, %r9, %r37; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 and.b32 %r146, %r2, 112; xor.b32 %r15, %r139, %r146; add.s32 %r65, %r142, %r15; selp.b32 %r66, 16, 0, %p10; // 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 [ %r65 + 0 ], [ %rd35 + 0 ], 0x10, %r66; // end inline asm add.s32 %r67, %r65, 4096; selp.b32 %r147, 16, 0, %p11; selp.b32 %r68, %r147, 0, %p7; // 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 [ %r67 + 0 ], [ %rd36 + 0 ], 0x10, %r68; // end inline asm add.s32 %r69, %r65, 8192; selp.b32 %r148, 16, 0, %p12; selp.b32 %r70, %r148, 0, %p7; // 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 [ %r69 + 0 ], [ %rd37 + 0 ], 0x10, %r70; // end inline asm add.s32 %r71, %r65, 12288; selp.b32 %r149, 16, 0, %p13; selp.b32 %r72, %r149, 0, %p7; // 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 [ %r71 + 0 ], [ %rd38 + 0 ], 0x10, %r72; // 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 %p14, %r37, 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, %r143; 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 %rd91, %r144, 2; add.s64 %rd42, %rd35, %rd91; add.s64 %rd43, %rd36, %rd91; add.s64 %rd44, %rd37, %rd91; add.s64 %rd45, %rd38, %rd91; .loc 1 15 37, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:15:37 add.s32 %r150, %r37, -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 %p15, %r3, %r150; setp.lt.s32 %p16, %r4, %r150; .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 %r73, %r13, 57344; selp.b32 %r151, 16, 0, %p15; selp.b32 %r74, %r151, 0, %p14; // 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 [ %r73 + 0 ], [ %rd39 + 0 ], 0x10, %r74; // end inline asm add.s32 %r75, %r13, 61440; selp.b32 %r152, 16, 0, %p16; selp.b32 %r76, %r152, 0, %p14; // 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 [ %r75 + 0 ], [ %rd40 + 0 ], 0x10, %r76; // 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 %r78, 4, 0, %p14; add.s32 %r77, %r142, 90116; // begin inline asm @%p5 cp.async.ca.shared.global [ %r77 + 0 ], [ %rd41 + 0 ], 0x4, %r78; // 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 %p17, %r6, %r150; setp.lt.s32 %p18, %r7, %r150; setp.lt.s32 %p19, %r8, %r150; setp.lt.s32 %p20, %r9, %r150; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 add.s32 %r79, %r65, 16384; selp.b32 %r153, 16, 0, %p17; selp.b32 %r80, %r153, 0, %p14; // 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 [ %r79 + 0 ], [ %rd42 + 0 ], 0x10, %r80; // end inline asm add.s32 %r81, %r65, 20480; selp.b32 %r154, 16, 0, %p18; selp.b32 %r82, %r154, 0, %p14; // 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 [ %r81 + 0 ], [ %rd43 + 0 ], 0x10, %r82; // end inline asm add.s32 %r83, %r65, 24576; selp.b32 %r155, 16, 0, %p19; selp.b32 %r84, %r155, 0, %p14; // 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 [ %r83 + 0 ], [ %rd44 + 0 ], 0x10, %r84; // end inline asm add.s32 %r85, %r65, 28672; selp.b32 %r156, 16, 0, %p20; selp.b32 %r86, %r156, 0, %p14; // 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 [ %r85 + 0 ], [ %rd45 + 0 ], 0x10, %r86; // 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 %p21, %r37, 128; .loc 1 34 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:34:9 add.s64 %rd46, %rd39, %rd16; add.s64 %rd47, %rd40, %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 %rd48, %rd34, 8; .loc 1 36 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:36:9 add.s64 %rd49, %rd42, %rd91; add.s64 %rd50, %rd43, %rd91; add.s64 %rd51, %rd44, %rd91; add.s64 %rd52, %rd45, %rd91; .loc 1 15 37, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:15:37 add.s32 %r157, %r37, -128; .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 %p22, %r3, %r157; setp.lt.s32 %p23, %r4, %r157; .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 %r87, %r13, 65536; selp.b32 %r158, 16, 0, %p22; selp.b32 %r88, %r158, 0, %p21; // 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 [ %r87 + 0 ], [ %rd46 + 0 ], 0x10, %r88; // end inline asm add.s32 %r89, %r13, 69632; selp.b32 %r159, 16, 0, %p23; selp.b32 %r90, %r159, 0, %p21; // 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 [ %r89 + 0 ], [ %rd47 + 0 ], 0x10, %r90; // 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 %r92, 4, 0, %p21; add.s32 %r91, %r142, 90120; // begin inline asm @%p5 cp.async.ca.shared.global [ %r91 + 0 ], [ %rd48 + 0 ], 0x4, %r92; // 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 %p24, %r6, %r157; setp.lt.s32 %p25, %r7, %r157; setp.lt.s32 %p26, %r8, %r157; setp.lt.s32 %p27, %r9, %r157; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 add.s32 %r93, %r65, 32768; selp.b32 %r160, 16, 0, %p24; selp.b32 %r94, %r160, 0, %p21; // 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 [ %r93 + 0 ], [ %rd49 + 0 ], 0x10, %r94; // end inline asm add.s32 %r95, %r65, 36864; selp.b32 %r161, 16, 0, %p25; selp.b32 %r96, %r161, 0, %p21; // 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 [ %r95 + 0 ], [ %rd50 + 0 ], 0x10, %r96; // end inline asm add.s32 %r97, %r65, 40960; selp.b32 %r162, 16, 0, %p26; selp.b32 %r98, %r162, 0, %p21; // 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 [ %r97 + 0 ], [ %rd51 + 0 ], 0x10, %r98; // end inline asm add.s32 %r99, %r65, 45056; selp.b32 %r163, 16, 0, %p27; selp.b32 %r100, %r163, 0, %p21; // 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 [ %r99 + 0 ], [ %rd52 + 0 ], 0x10, %r100; // end inline asm cp.async.commit_group; mov.b32 %r1003, 0f00000000; mov.b32 %r1001, 0f3F800000; mov.b32 %r1002, %r1001; mov.b32 %r1004, %r1003; mov.b32 %r1005, %r1003; mov.b32 %r1006, %r1003; mov.b32 %r1007, %r1003; mov.b32 %r1008, %r1003; mov.b32 %r1009, %r1003; mov.b32 %r1010, %r1003; mov.b32 %r1011, %r1003; mov.b32 %r1012, %r1003; mov.b32 %r1013, %r1003; mov.b32 %r1014, %r1003; mov.b32 %r1015, %r1003; mov.b32 %r1016, %r1003; mov.b32 %r1017, %r1003; mov.b32 %r1018, %r1003; mov.b32 %r1019, %r1003; mov.b32 %r1020, %r1003; mov.b32 %r1021, %r1003; mov.b32 %r1022, %r1003; mov.b32 %r1023, %r1003; mov.b32 %r1024, %r1003; mov.b32 %r1025, %r1003; mov.b32 %r1026, %r1003; mov.b32 %r1027, %r1003; mov.b32 %r1028, %r1003; mov.b32 %r1029, %r1003; mov.b32 %r1030, %r1003; mov.b32 %r1031, %r1003; mov.b32 %r1032, %r1003; mov.b32 %r1033, %r1003; mov.b32 %r1034, %r1003; mov.b32 %r1035, %r1003; mov.b32 %r1036, %r1003; mov.b32 %r1037, %r1003; mov.b32 %r1038, %r1003; mov.b32 %r1039, %r1003; mov.b32 %r1040, %r1003; mov.b32 %r1041, %r1003; mov.b32 %r1042, %r1003; mov.b32 %r1043, %r1003; mov.b32 %r1044, %r1003; mov.b32 %r1045, %r1003; mov.b32 %r1046, %r1003; mov.b32 %r1047, %r1003; mov.b32 %r1048, %r1003; mov.b32 %r1049, %r1003; mov.b32 %r1050, %r1003; mov.b32 %r1051, %r1003; mov.b32 %r1052, %r1003; mov.b32 %r1053, %r1003; mov.b32 %r1054, %r1003; mov.b32 %r1055, %r1003; mov.b32 %r1056, %r1003; mov.b32 %r1057, %r1003; mov.b32 %r1058, %r1003; mov.b32 %r1059, %r1003; mov.b32 %r1060, %r1003; mov.b32 %r1061, %r1003; mov.b32 %r1062, %r1003; mov.b32 %r1063, %r1003; mov.b32 %r1064, %r1003; mov.b32 %r1065, %r1003; mov.b32 %r1066, %r1003; .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_3; // %bb.1: // %.lr.ph .loc 1 0 5 // attn_qk_int8_per_block.py:0:5 cvt.s64.s32 %rd11, %r132; cvt.s64.s32 %rd12, %r134; cvt.s64.s32 %rd13, %r136; cvt.s64.s32 %rd14, %r137; cvt.s64.s32 %rd17, %r144; cvt.u32.u64 %r164, %rd4; add.s32 %r16, %r37, -192; shl.b32 %r165, %r11, 7; shl.b32 %r166, %r2, 6; and.b32 %r167, %r166, 14336; and.b32 %r168, %r2, 16; or.b32 %r169, %r165, %r167; xor.b32 %r170, %r164, %r168; or.b32 %r171, %r169, %r170; add.s32 %r172, %r142, %r171; add.s32 %r17, %r172, 73728; xor.b32 %r173, %r171, 32; add.s32 %r174, %r142, %r173; add.s32 %r18, %r174, 73728; xor.b32 %r175, %r171, 64; add.s32 %r176, %r142, %r175; add.s32 %r19, %r176, 73728; xor.b32 %r177, %r171, 96; add.s32 %r178, %r142, %r177; add.s32 %r20, %r178, 73728; and.b32 %r179, %r12, 48; mul.lo.s32 %r180, %r10, 144; xor.b32 %r21, %r180, %r179; xor.b32 %r22, %r21, 64; shl.b32 %r181, %r2, 8; and.b32 %r182, %r181, 7936; or.b32 %r23, %r182, %r164; xor.b32 %r24, %r23, 16; xor.b32 %r25, %r23, 32; xor.b32 %r26, %r23, 48; xor.b32 %r27, %r23, 64; xor.b32 %r28, %r23, 80; xor.b32 %r29, %r23, 96; xor.b32 %r30, %r23, 112; .loc 1 13 5 // attn_qk_int8_per_block.py:13:5 @[ attn_qk_int8_per_block.py:75:16 ] add.s64 %rd92, %rd9, %rd10; add.s64 %rd93, %rd92, %rd14; shl.b64 %rd18, %rd93, 1; mul.lo.s64 %rd94, %rd17, 6; mad.wide.u32 %rd95, %r11, 16, %rd94; add.s64 %rd136, %rd24, %rd95; shl.b64 %rd19, %rd17, 1; add.s64 %rd96, %rd92, %rd13; shl.b64 %rd20, %rd96, 1; add.s64 %rd97, %rd92, %rd12; shl.b64 %rd21, %rd97, 1; add.s64 %rd98, %rd92, %rd11; shl.b64 %rd22, %rd98, 1; add.s64 %rd99, %rd6, %rd5; add.s64 %rd100, %rd99, %rd4; mad.lo.s64 %rd101, %rd16, 3, %rd100; add.s64 %rd135, %rd23, %rd101; add.s64 %rd102, %rd81, %rd25; add.s64 %rd134, %rd102, 12; mov.b32 %r1003, 0f00000000; mov.b32 %r183, 0; mov.b32 %r1001, 0f3F800000; mov.b32 %r998, 0fFF800000; mov.b32 %r997, 2; mov.b32 %r996, -1; mov.b32 %r995, %r16; mov.b32 %r999, %r998; mov.b32 %r1002, %r1001; mov.b32 %r1000, %r183; mov.b32 %r1004, %r1003; mov.b32 %r1005, %r1003; mov.b32 %r1006, %r1003; mov.b32 %r1007, %r1003; mov.b32 %r1008, %r1003; mov.b32 %r1009, %r1003; mov.b32 %r1010, %r1003; mov.b32 %r1011, %r1003; mov.b32 %r1012, %r1003; mov.b32 %r1013, %r1003; mov.b32 %r1014, %r1003; mov.b32 %r1015, %r1003; mov.b32 %r1016, %r1003; mov.b32 %r1017, %r1003; mov.b32 %r1018, %r1003; mov.b32 %r1019, %r1003; mov.b32 %r1020, %r1003; mov.b32 %r1021, %r1003; mov.b32 %r1022, %r1003; mov.b32 %r1023, %r1003; mov.b32 %r1024, %r1003; mov.b32 %r1025, %r1003; mov.b32 %r1026, %r1003; mov.b32 %r1027, %r1003; mov.b32 %r1028, %r1003; mov.b32 %r1029, %r1003; mov.b32 %r1030, %r1003; mov.b32 %r1031, %r1003; mov.b32 %r1032, %r1003; mov.b32 %r1033, %r1003; mov.b32 %r1034, %r1003; mov.b32 %r1035, %r1003; mov.b32 %r1036, %r1003; mov.b32 %r1037, %r1003; mov.b32 %r1038, %r1003; mov.b32 %r1039, %r1003; mov.b32 %r1040, %r1003; mov.b32 %r1041, %r1003; mov.b32 %r1042, %r1003; mov.b32 %r1043, %r1003; mov.b32 %r1044, %r1003; mov.b32 %r1045, %r1003; mov.b32 %r1046, %r1003; mov.b32 %r1047, %r1003; mov.b32 %r1048, %r1003; mov.b32 %r1049, %r1003; mov.b32 %r1050, %r1003; mov.b32 %r1051, %r1003; mov.b32 %r1052, %r1003; mov.b32 %r1053, %r1003; mov.b32 %r1054, %r1003; mov.b32 %r1055, %r1003; mov.b32 %r1056, %r1003; mov.b32 %r1057, %r1003; mov.b32 %r1058, %r1003; mov.b32 %r1059, %r1003; mov.b32 %r1060, %r1003; mov.b32 %r1061, %r1003; mov.b32 %r1062, %r1003; mov.b32 %r1063, %r1003; mov.b32 %r1064, %r1003; mov.b32 %r1065, %r1003; mov.b32 %r1066, %r1003; $L__BB0_2: // %__nv_exp2f.exit // =>This Inner Loop Header: Depth=1 setp.lt.s32 %p28, %r1000, %r16; add.s32 %r486, %r996, 1; setp.gt.s32 %p29, %r486, 2; selp.b32 %r996, 0, %r486, %p29; .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 6; bar.sync 0; shl.b32 %r487, %r996, 13; add.s32 %r488, %r142, %r487; .loc 1 17 19, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:17:19 shl.b32 %r489, %r996, 2; add.s32 %r490, %r142, %r489; ld.shared.b32 %r491, [%r490+90112]; $L__tmp4: .loc 1 73 9 // attn_qk_int8_per_block.py:73:9 ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r184, %r185, %r186, %r187}, [%r17]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r208, %r209, %r210, %r211}, [%r18]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r256, %r257, %r258, %r259}, [%r19]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r276, %r277, %r278, %r279}, [%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 %r492, %r488, %r21; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r188, %r189, %r212, %r213}, [%r492+49152]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r190, %r191, %r218, %r219}, [%r492+50176]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r192, %r193, %r224, %r225}, [%r492+51200]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r194, %r195, %r230, %r231}, [%r492+52224]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r196, %r197, %r236, %r237}, [%r492+53248]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r198, %r199, %r242, %r243}, [%r492+54272]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r200, %r201, %r248, %r249}, [%r492+55296]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r202, %r203, %r254, %r255}, [%r492+56320]; add.s32 %r493, %r488, %r22; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r260, %r261, %r280, %r281}, [%r493+49152]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r262, %r263, %r282, %r283}, [%r493+50176]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r264, %r265, %r284, %r285}, [%r493+51200]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r266, %r267, %r286, %r287}, [%r493+52224]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r268, %r269, %r288, %r289}, [%r493+53248]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r270, %r271, %r290, %r291}, [%r493+54272]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r272, %r273, %r292, %r293}, [%r493+55296]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r274, %r275, %r294, %r295}, [%r493+56320]; .loc 1 18 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:18:14 mov.b32 %r204, %r183; mov.b32 %r205, %r183; mov.b32 %r206, %r183; mov.b32 %r207, %r183; // 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 { %r204, %r205, %r206, %r207 }, { %r184, %r185, %r186, %r187 }, { %r188, %r189 }, { %r204, %r205, %r206, %r207 }; // end inline asm mov.b32 %r214, %r183; mov.b32 %r215, %r183; mov.b32 %r216, %r183; mov.b32 %r217, %r183; // 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 { %r214, %r215, %r216, %r217 }, { %r184, %r185, %r186, %r187 }, { %r190, %r191 }, { %r214, %r215, %r216, %r217 }; // end inline asm mov.b32 %r220, %r183; mov.b32 %r221, %r183; mov.b32 %r222, %r183; mov.b32 %r223, %r183; // 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 { %r220, %r221, %r222, %r223 }, { %r184, %r185, %r186, %r187 }, { %r192, %r193 }, { %r220, %r221, %r222, %r223 }; // end inline asm mov.b32 %r226, %r183; mov.b32 %r227, %r183; mov.b32 %r228, %r183; mov.b32 %r229, %r183; // 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 { %r226, %r227, %r228, %r229 }, { %r184, %r185, %r186, %r187 }, { %r194, %r195 }, { %r226, %r227, %r228, %r229 }; // end inline asm mov.b32 %r232, %r183; mov.b32 %r233, %r183; mov.b32 %r234, %r183; mov.b32 %r235, %r183; // 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 { %r232, %r233, %r234, %r235 }, { %r184, %r185, %r186, %r187 }, { %r196, %r197 }, { %r232, %r233, %r234, %r235 }; // end inline asm mov.b32 %r238, %r183; mov.b32 %r239, %r183; mov.b32 %r240, %r183; mov.b32 %r241, %r183; // 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 { %r238, %r239, %r240, %r241 }, { %r184, %r185, %r186, %r187 }, { %r198, %r199 }, { %r238, %r239, %r240, %r241 }; // end inline asm mov.b32 %r244, %r183; mov.b32 %r245, %r183; mov.b32 %r246, %r183; mov.b32 %r247, %r183; // 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 }, { %r184, %r185, %r186, %r187 }, { %r200, %r201 }, { %r244, %r245, %r246, %r247 }; // end inline asm mov.b32 %r250, %r183; mov.b32 %r251, %r183; mov.b32 %r252, %r183; mov.b32 %r253, %r183; // 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 { %r250, %r251, %r252, %r253 }, { %r184, %r185, %r186, %r187 }, { %r202, %r203 }, { %r250, %r251, %r252, %r253 }; // 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 { %r204, %r205, %r206, %r207 }, { %r208, %r209, %r210, %r211 }, { %r212, %r213 }, { %r204, %r205, %r206, %r207 }; // 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 { %r214, %r215, %r216, %r217 }, { %r208, %r209, %r210, %r211 }, { %r218, %r219 }, { %r214, %r215, %r216, %r217 }; // 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 { %r220, %r221, %r222, %r223 }, { %r208, %r209, %r210, %r211 }, { %r224, %r225 }, { %r220, %r221, %r222, %r223 }; // 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 { %r226, %r227, %r228, %r229 }, { %r208, %r209, %r210, %r211 }, { %r230, %r231 }, { %r226, %r227, %r228, %r229 }; // 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 { %r232, %r233, %r234, %r235 }, { %r208, %r209, %r210, %r211 }, { %r236, %r237 }, { %r232, %r233, %r234, %r235 }; // 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 { %r238, %r239, %r240, %r241 }, { %r208, %r209, %r210, %r211 }, { %r242, %r243 }, { %r238, %r239, %r240, %r241 }; // 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 }, { %r208, %r209, %r210, %r211 }, { %r248, %r249 }, { %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 { %r250, %r251, %r252, %r253 }, { %r208, %r209, %r210, %r211 }, { %r254, %r255 }, { %r250, %r251, %r252, %r253 }; // 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 { %r204, %r205, %r206, %r207 }, { %r256, %r257, %r258, %r259 }, { %r260, %r261 }, { %r204, %r205, %r206, %r207 }; // 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 { %r214, %r215, %r216, %r217 }, { %r256, %r257, %r258, %r259 }, { %r262, %r263 }, { %r214, %r215, %r216, %r217 }; // 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 { %r220, %r221, %r222, %r223 }, { %r256, %r257, %r258, %r259 }, { %r264, %r265 }, { %r220, %r221, %r222, %r223 }; // 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 { %r226, %r227, %r228, %r229 }, { %r256, %r257, %r258, %r259 }, { %r266, %r267 }, { %r226, %r227, %r228, %r229 }; // 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 { %r232, %r233, %r234, %r235 }, { %r256, %r257, %r258, %r259 }, { %r268, %r269 }, { %r232, %r233, %r234, %r235 }; // 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 { %r238, %r239, %r240, %r241 }, { %r256, %r257, %r258, %r259 }, { %r270, %r271 }, { %r238, %r239, %r240, %r241 }; // 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 }, { %r256, %r257, %r258, %r259 }, { %r272, %r273 }, { %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 { %r250, %r251, %r252, %r253 }, { %r256, %r257, %r258, %r259 }, { %r274, %r275 }, { %r250, %r251, %r252, %r253 }; // 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 { %r204, %r205, %r206, %r207 }, { %r276, %r277, %r278, %r279 }, { %r280, %r281 }, { %r204, %r205, %r206, %r207 }; // 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 { %r214, %r215, %r216, %r217 }, { %r276, %r277, %r278, %r279 }, { %r282, %r283 }, { %r214, %r215, %r216, %r217 }; // 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 { %r220, %r221, %r222, %r223 }, { %r276, %r277, %r278, %r279 }, { %r284, %r285 }, { %r220, %r221, %r222, %r223 }; // 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 { %r226, %r227, %r228, %r229 }, { %r276, %r277, %r278, %r279 }, { %r286, %r287 }, { %r226, %r227, %r228, %r229 }; // 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 { %r232, %r233, %r234, %r235 }, { %r276, %r277, %r278, %r279 }, { %r288, %r289 }, { %r232, %r233, %r234, %r235 }; // 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 { %r238, %r239, %r240, %r241 }, { %r276, %r277, %r278, %r279 }, { %r290, %r291 }, { %r238, %r239, %r240, %r241 }; // 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 }, { %r276, %r277, %r278, %r279 }, { %r292, %r293 }, { %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 { %r250, %r251, %r252, %r253 }, { %r276, %r277, %r278, %r279 }, { %r294, %r295 }, { %r250, %r251, %r252, %r253 }; // end inline asm cvt.rn.f32.s32 %r494, %r251; cvt.rn.f32.s32 %r495, %r250; cvt.rn.f32.s32 %r496, %r245; cvt.rn.f32.s32 %r497, %r244; cvt.rn.f32.s32 %r498, %r239; cvt.rn.f32.s32 %r499, %r238; cvt.rn.f32.s32 %r500, %r233; cvt.rn.f32.s32 %r501, %r232; cvt.rn.f32.s32 %r502, %r227; cvt.rn.f32.s32 %r503, %r226; cvt.rn.f32.s32 %r504, %r221; cvt.rn.f32.s32 %r505, %r220; cvt.rn.f32.s32 %r506, %r215; cvt.rn.f32.s32 %r507, %r214; cvt.rn.f32.s32 %r508, %r205; cvt.rn.f32.s32 %r509, %r204; mul.f32 %r510, %r58, %r509; mul.f32 %r511, %r58, %r508; mul.f32 %r512, %r58, %r507; mul.f32 %r513, %r58, %r506; mul.f32 %r514, %r58, %r505; mul.f32 %r515, %r58, %r504; mul.f32 %r516, %r58, %r503; mul.f32 %r517, %r58, %r502; mul.f32 %r518, %r58, %r501; mul.f32 %r519, %r58, %r500; mul.f32 %r520, %r58, %r499; mul.f32 %r521, %r58, %r498; mul.f32 %r522, %r58, %r497; mul.f32 %r523, %r58, %r496; mul.f32 %r524, %r58, %r495; mul.f32 %r525, %r58, %r494; mul.f32 %r526, %r491, %r525; mul.f32 %r527, %r491, %r524; mul.f32 %r528, %r491, %r523; mul.f32 %r529, %r491, %r522; mul.f32 %r530, %r491, %r521; mul.f32 %r531, %r491, %r520; mul.f32 %r532, %r491, %r519; mul.f32 %r533, %r491, %r518; mul.f32 %r534, %r491, %r517; mul.f32 %r535, %r491, %r516; mul.f32 %r536, %r491, %r515; mul.f32 %r537, %r491, %r514; mul.f32 %r538, %r491, %r513; mul.f32 %r539, %r491, %r512; mul.f32 %r540, %r491, %r511; mul.f32 %r541, %r491, %r510; cvt.rn.f32.s32 %r542, %r253; cvt.rn.f32.s32 %r543, %r252; cvt.rn.f32.s32 %r544, %r247; cvt.rn.f32.s32 %r545, %r246; cvt.rn.f32.s32 %r546, %r241; cvt.rn.f32.s32 %r547, %r240; cvt.rn.f32.s32 %r548, %r235; cvt.rn.f32.s32 %r549, %r234; cvt.rn.f32.s32 %r550, %r229; cvt.rn.f32.s32 %r551, %r228; cvt.rn.f32.s32 %r552, %r223; cvt.rn.f32.s32 %r553, %r222; cvt.rn.f32.s32 %r554, %r217; cvt.rn.f32.s32 %r555, %r216; cvt.rn.f32.s32 %r556, %r207; cvt.rn.f32.s32 %r557, %r206; mul.f32 %r558, %r58, %r557; mul.f32 %r559, %r58, %r556; mul.f32 %r560, %r58, %r555; mul.f32 %r561, %r58, %r554; mul.f32 %r562, %r58, %r553; mul.f32 %r563, %r58, %r552; mul.f32 %r564, %r58, %r551; mul.f32 %r565, %r58, %r550; mul.f32 %r566, %r58, %r549; mul.f32 %r567, %r58, %r548; mul.f32 %r568, %r58, %r547; mul.f32 %r569, %r58, %r546; mul.f32 %r570, %r58, %r545; mul.f32 %r571, %r58, %r544; mul.f32 %r572, %r58, %r543; mul.f32 %r573, %r58, %r542; mul.f32 %r574, %r491, %r573; mul.f32 %r575, %r491, %r572; mul.f32 %r576, %r491, %r571; mul.f32 %r577, %r491, %r570; mul.f32 %r578, %r491, %r569; mul.f32 %r579, %r491, %r568; mul.f32 %r580, %r491, %r567; mul.f32 %r581, %r491, %r566; mul.f32 %r582, %r491, %r565; mul.f32 %r583, %r491, %r564; mul.f32 %r584, %r491, %r563; mul.f32 %r585, %r491, %r562; mul.f32 %r586, %r491, %r561; mul.f32 %r587, %r491, %r560; mul.f32 %r588, %r491, %r559; mul.f32 %r589, %r491, %r558; $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 %r590, %r541, %r540; max.f32 %r591, %r539, %r538; max.f32 %r592, %r537, %r536; max.f32 %r593, %r535, %r534; max.f32 %r594, %r533, %r532; max.f32 %r595, %r531, %r530; max.f32 %r596, %r529, %r528; max.f32 %r597, %r527, %r526; max.f32 %r598, %r590, %r591; max.f32 %r599, %r592, %r593; max.f32 %r600, %r594, %r595; max.f32 %r601, %r596, %r597; max.f32 %r602, %r598, %r599; max.f32 %r603, %r600, %r601; max.f32 %r604, %r602, %r603; max.f32 %r605, %r589, %r588; max.f32 %r606, %r587, %r586; max.f32 %r607, %r585, %r584; max.f32 %r608, %r583, %r582; max.f32 %r609, %r581, %r580; max.f32 %r610, %r579, %r578; max.f32 %r611, %r577, %r576; max.f32 %r612, %r575, %r574; max.f32 %r613, %r605, %r606; max.f32 %r614, %r607, %r608; max.f32 %r615, %r609, %r610; max.f32 %r616, %r611, %r612; max.f32 %r617, %r613, %r614; max.f32 %r618, %r615, %r616; max.f32 %r619, %r617, %r618; $L__tmp7: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r620, %r604, 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 %r621, %r604, %r620; $L__tmp9: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r622, %r621, 1, 31, -1; $L__tmp10: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r623, %r621, %r622; $L__tmp11: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r624, %r619, 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 %r625, %r619, %r624; $L__tmp13: .loc 2 191 16, function_name $L__info_string0, inlined_at 1 19 32 // standard.py:191:16 shfl.sync.bfly.b32 %r626, %r625, 1, 31, -1; $L__tmp14: .loc 2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12 max.f32 %r627, %r625, %r626; $L__tmp15: .loc 1 19 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:19:16 max.f32 %r31, %r998, %r623; max.f32 %r32, %r999, %r627; .loc 1 20 14, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:20:14 neg.f32 %r628, %r31; fma.rn.f32 %r629, %r491, %r510, %r628; fma.rn.f32 %r630, %r491, %r511, %r628; neg.f32 %r631, %r32; fma.rn.f32 %r632, %r491, %r558, %r631; fma.rn.f32 %r633, %r491, %r559, %r631; fma.rn.f32 %r634, %r491, %r512, %r628; fma.rn.f32 %r635, %r491, %r513, %r628; fma.rn.f32 %r636, %r491, %r560, %r631; fma.rn.f32 %r637, %r491, %r561, %r631; fma.rn.f32 %r638, %r491, %r514, %r628; fma.rn.f32 %r639, %r491, %r515, %r628; fma.rn.f32 %r640, %r491, %r562, %r631; fma.rn.f32 %r641, %r491, %r563, %r631; fma.rn.f32 %r642, %r491, %r516, %r628; fma.rn.f32 %r643, %r491, %r517, %r628; fma.rn.f32 %r644, %r491, %r564, %r631; fma.rn.f32 %r645, %r491, %r565, %r631; fma.rn.f32 %r646, %r491, %r518, %r628; fma.rn.f32 %r647, %r491, %r519, %r628; fma.rn.f32 %r648, %r491, %r566, %r631; fma.rn.f32 %r649, %r491, %r567, %r631; fma.rn.f32 %r650, %r491, %r520, %r628; fma.rn.f32 %r651, %r491, %r521, %r628; fma.rn.f32 %r652, %r491, %r568, %r631; fma.rn.f32 %r653, %r491, %r569, %r631; fma.rn.f32 %r654, %r491, %r522, %r628; fma.rn.f32 %r655, %r491, %r523, %r628; fma.rn.f32 %r656, %r491, %r570, %r631; fma.rn.f32 %r657, %r491, %r571, %r631; fma.rn.f32 %r658, %r491, %r524, %r628; fma.rn.f32 %r659, %r491, %r525, %r628; fma.rn.f32 %r660, %r491, %r572, %r631; fma.rn.f32 %r661, %r491, %r573, %r631; .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 %r662, %r629; ex2.approx.ftz.f32 %r663, %r630; ex2.approx.ftz.f32 %r664, %r632; ex2.approx.ftz.f32 %r665, %r633; ex2.approx.ftz.f32 %r666, %r634; ex2.approx.ftz.f32 %r667, %r635; ex2.approx.ftz.f32 %r668, %r636; ex2.approx.ftz.f32 %r669, %r637; ex2.approx.ftz.f32 %r670, %r638; ex2.approx.ftz.f32 %r671, %r639; ex2.approx.ftz.f32 %r672, %r640; ex2.approx.ftz.f32 %r673, %r641; ex2.approx.ftz.f32 %r674, %r642; ex2.approx.ftz.f32 %r675, %r643; ex2.approx.ftz.f32 %r676, %r644; ex2.approx.ftz.f32 %r677, %r645; ex2.approx.ftz.f32 %r678, %r646; ex2.approx.ftz.f32 %r679, %r647; ex2.approx.ftz.f32 %r680, %r648; ex2.approx.ftz.f32 %r681, %r649; ex2.approx.ftz.f32 %r682, %r650; ex2.approx.ftz.f32 %r683, %r651; ex2.approx.ftz.f32 %r684, %r652; ex2.approx.ftz.f32 %r685, %r653; ex2.approx.ftz.f32 %r686, %r654; ex2.approx.ftz.f32 %r687, %r655; ex2.approx.ftz.f32 %r688, %r656; ex2.approx.ftz.f32 %r689, %r657; ex2.approx.ftz.f32 %r690, %r658; ex2.approx.ftz.f32 %r691, %r659; ex2.approx.ftz.f32 %r692, %r660; ex2.approx.ftz.f32 %r693, %r661; $L__tmp16: .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 %r694, %r662, %r663; add.f32 %r695, %r666, %r667; add.f32 %r696, %r670, %r671; add.f32 %r697, %r674, %r675; add.f32 %r698, %r678, %r679; add.f32 %r699, %r682, %r683; add.f32 %r700, %r686, %r687; add.f32 %r701, %r690, %r691; add.f32 %r702, %r694, %r695; add.f32 %r703, %r696, %r697; add.f32 %r704, %r698, %r699; add.f32 %r705, %r700, %r701; add.f32 %r706, %r702, %r703; add.f32 %r707, %r704, %r705; add.f32 %r708, %r706, %r707; add.f32 %r709, %r664, %r665; add.f32 %r710, %r668, %r669; add.f32 %r711, %r672, %r673; add.f32 %r712, %r676, %r677; add.f32 %r713, %r680, %r681; add.f32 %r714, %r684, %r685; add.f32 %r715, %r688, %r689; add.f32 %r716, %r692, %r693; add.f32 %r717, %r709, %r710; add.f32 %r718, %r711, %r712; add.f32 %r719, %r713, %r714; add.f32 %r720, %r715, %r716; add.f32 %r721, %r717, %r718; add.f32 %r722, %r719, %r720; add.f32 %r723, %r721, %r722; $L__tmp17: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r724, %r708, 2, 31, -1; $L__tmp18: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r725, %r708, %r724; $L__tmp19: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r726, %r725, 1, 31, -1; $L__tmp20: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r727, %r725, %r726; $L__tmp21: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r728, %r723, 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 %r729, %r723, %r728; $L__tmp23: .loc 2 293 12, function_name $L__info_string0, inlined_at 1 22 16 // standard.py:293:12 shfl.sync.bfly.b32 %r730, %r729, 1, 31, -1; $L__tmp24: .loc 2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12 add.f32 %r731, %r729, %r730; $L__tmp25: .loc 1 24 30, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:24:30 sub.f32 %r732, %r998, %r31; sub.f32 %r733, %r999, %r32; .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 %r734, %r732; ex2.approx.ftz.f32 %r735, %r733; .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 %r1001, %r1001, %r734, %r727; fma.rn.f32 %r1002, %r1002, %r735, %r731; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 add.s32 %r736, %r488, %r487; add.s32 %r737, %r736, %r23; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r296, %r297, %r334, %r335}, [%r737]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r400, %r401, %r436, %r437}, [%r737+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r316, %r317, %r370, %r371}, [%r737+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r420, %r421, %r456, %r457}, [%r737+8320]; add.s32 %r738, %r736, %r24; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r302, %r303, %r342, %r343}, [%r738]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r406, %r407, %r442, %r443}, [%r738+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r318, %r319, %r374, %r375}, [%r738+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r422, %r423, %r458, %r459}, [%r738+8320]; add.s32 %r739, %r736, %r25; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r304, %r305, %r346, %r347}, [%r739]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r408, %r409, %r444, %r445}, [%r739+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r320, %r321, %r378, %r379}, [%r739+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r424, %r425, %r460, %r461}, [%r739+8320]; add.s32 %r740, %r736, %r26; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r306, %r307, %r350, %r351}, [%r740]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r410, %r411, %r446, %r447}, [%r740+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r322, %r323, %r382, %r383}, [%r740+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r426, %r427, %r462, %r463}, [%r740+8320]; add.s32 %r741, %r736, %r27; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r308, %r309, %r354, %r355}, [%r741]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r412, %r413, %r448, %r449}, [%r741+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r324, %r325, %r386, %r387}, [%r741+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r428, %r429, %r464, %r465}, [%r741+8320]; add.s32 %r742, %r736, %r28; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r310, %r311, %r358, %r359}, [%r742]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r414, %r415, %r450, %r451}, [%r742+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r326, %r327, %r390, %r391}, [%r742+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r430, %r431, %r466, %r467}, [%r742+8320]; add.s32 %r743, %r736, %r29; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r312, %r313, %r362, %r363}, [%r743]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r416, %r417, %r452, %r453}, [%r743+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r328, %r329, %r394, %r395}, [%r743+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r432, %r433, %r468, %r469}, [%r743+8320]; add.s32 %r744, %r736, %r30; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r314, %r315, %r366, %r367}, [%r744]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r418, %r419, %r454, %r455}, [%r744+8192]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r330, %r331, %r398, %r399}, [%r744+128]; ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r434, %r435, %r470, %r471}, [%r744+8320]; .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 %r298, %r663, %r662; cvt.rn.f16x2.f32 %r299, %r665, %r664; cvt.rn.f16x2.f32 %r300, %r667, %r666; cvt.rn.f16x2.f32 %r301, %r669, %r668; cvt.rn.f16x2.f32 %r338, %r671, %r670; cvt.rn.f16x2.f32 %r339, %r673, %r672; cvt.rn.f16x2.f32 %r340, %r675, %r674; cvt.rn.f16x2.f32 %r341, %r677, %r676; cvt.rn.f16x2.f32 %r402, %r679, %r678; cvt.rn.f16x2.f32 %r403, %r681, %r680; cvt.rn.f16x2.f32 %r404, %r683, %r682; cvt.rn.f16x2.f32 %r405, %r685, %r684; cvt.rn.f16x2.f32 %r438, %r687, %r686; cvt.rn.f16x2.f32 %r439, %r689, %r688; cvt.rn.f16x2.f32 %r440, %r691, %r690; cvt.rn.f16x2.f32 %r441, %r693, %r692; .loc 1 32 16, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:32:16 mov.b32 %r333, %r183; mov.b32 %r332, %r183; // 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 }, { %r298, %r299, %r300, %r301 }, { %r296, %r297 }, { %r332, %r333 }; // end inline asm mov.b32 %r336, %r183; mov.b32 %r337, %r183; // 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 { %r336, %r337 }, { %r298, %r299, %r300, %r301 }, { %r302, %r303 }, { %r336, %r337 }; // end inline asm mov.b32 %r344, %r183; mov.b32 %r345, %r183; // 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 }, { %r298, %r299, %r300, %r301 }, { %r304, %r305 }, { %r344, %r345 }; // end inline asm mov.b32 %r348, %r183; mov.b32 %r349, %r183; // 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 }, { %r298, %r299, %r300, %r301 }, { %r306, %r307 }, { %r348, %r349 }; // end inline asm mov.b32 %r352, %r183; mov.b32 %r353, %r183; // 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 { %r352, %r353 }, { %r298, %r299, %r300, %r301 }, { %r308, %r309 }, { %r352, %r353 }; // end inline asm mov.b32 %r356, %r183; mov.b32 %r357, %r183; // 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 { %r356, %r357 }, { %r298, %r299, %r300, %r301 }, { %r310, %r311 }, { %r356, %r357 }; // end inline asm mov.b32 %r360, %r183; mov.b32 %r361, %r183; // 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 { %r360, %r361 }, { %r298, %r299, %r300, %r301 }, { %r312, %r313 }, { %r360, %r361 }; // end inline asm mov.b32 %r364, %r183; mov.b32 %r365, %r183; // 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 { %r364, %r365 }, { %r298, %r299, %r300, %r301 }, { %r314, %r315 }, { %r364, %r365 }; // end inline asm mov.b32 %r368, %r183; mov.b32 %r369, %r183; // 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 { %r368, %r369 }, { %r298, %r299, %r300, %r301 }, { %r316, %r317 }, { %r368, %r369 }; // end inline asm mov.b32 %r372, %r183; mov.b32 %r373, %r183; // 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 { %r372, %r373 }, { %r298, %r299, %r300, %r301 }, { %r318, %r319 }, { %r372, %r373 }; // end inline asm mov.b32 %r376, %r183; mov.b32 %r377, %r183; // 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 { %r376, %r377 }, { %r298, %r299, %r300, %r301 }, { %r320, %r321 }, { %r376, %r377 }; // end inline asm mov.b32 %r380, %r183; mov.b32 %r381, %r183; // 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 { %r380, %r381 }, { %r298, %r299, %r300, %r301 }, { %r322, %r323 }, { %r380, %r381 }; // end inline asm mov.b32 %r384, %r183; mov.b32 %r385, %r183; // 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 { %r384, %r385 }, { %r298, %r299, %r300, %r301 }, { %r324, %r325 }, { %r384, %r385 }; // end inline asm mov.b32 %r388, %r183; mov.b32 %r389, %r183; // 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 { %r388, %r389 }, { %r298, %r299, %r300, %r301 }, { %r326, %r327 }, { %r388, %r389 }; // end inline asm mov.b32 %r392, %r183; mov.b32 %r393, %r183; // 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 { %r392, %r393 }, { %r298, %r299, %r300, %r301 }, { %r328, %r329 }, { %r392, %r393 }; // end inline asm mov.b32 %r396, %r183; mov.b32 %r397, %r183; // 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 { %r396, %r397 }, { %r298, %r299, %r300, %r301 }, { %r330, %r331 }, { %r396, %r397 }; // 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 }, { %r338, %r339, %r340, %r341 }, { %r334, %r335 }, { %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 { %r336, %r337 }, { %r338, %r339, %r340, %r341 }, { %r342, %r343 }, { %r336, %r337 }; // 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 }, { %r338, %r339, %r340, %r341 }, { %r346, %r347 }, { %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 { %r348, %r349 }, { %r338, %r339, %r340, %r341 }, { %r350, %r351 }, { %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 { %r352, %r353 }, { %r338, %r339, %r340, %r341 }, { %r354, %r355 }, { %r352, %r353 }; // 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 { %r356, %r357 }, { %r338, %r339, %r340, %r341 }, { %r358, %r359 }, { %r356, %r357 }; // 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 { %r360, %r361 }, { %r338, %r339, %r340, %r341 }, { %r362, %r363 }, { %r360, %r361 }; // 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 { %r364, %r365 }, { %r338, %r339, %r340, %r341 }, { %r366, %r367 }, { %r364, %r365 }; // 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 { %r368, %r369 }, { %r338, %r339, %r340, %r341 }, { %r370, %r371 }, { %r368, %r369 }; // 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 { %r372, %r373 }, { %r338, %r339, %r340, %r341 }, { %r374, %r375 }, { %r372, %r373 }; // 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 { %r376, %r377 }, { %r338, %r339, %r340, %r341 }, { %r378, %r379 }, { %r376, %r377 }; // 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 { %r380, %r381 }, { %r338, %r339, %r340, %r341 }, { %r382, %r383 }, { %r380, %r381 }; // 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 { %r384, %r385 }, { %r338, %r339, %r340, %r341 }, { %r386, %r387 }, { %r384, %r385 }; // 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 { %r388, %r389 }, { %r338, %r339, %r340, %r341 }, { %r390, %r391 }, { %r388, %r389 }; // 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 { %r392, %r393 }, { %r338, %r339, %r340, %r341 }, { %r394, %r395 }, { %r392, %r393 }; // 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 { %r396, %r397 }, { %r338, %r339, %r340, %r341 }, { %r398, %r399 }, { %r396, %r397 }; // 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 }, { %r402, %r403, %r404, %r405 }, { %r400, %r401 }, { %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 { %r336, %r337 }, { %r402, %r403, %r404, %r405 }, { %r406, %r407 }, { %r336, %r337 }; // 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 }, { %r402, %r403, %r404, %r405 }, { %r408, %r409 }, { %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 { %r348, %r349 }, { %r402, %r403, %r404, %r405 }, { %r410, %r411 }, { %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 { %r352, %r353 }, { %r402, %r403, %r404, %r405 }, { %r412, %r413 }, { %r352, %r353 }; // 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 { %r356, %r357 }, { %r402, %r403, %r404, %r405 }, { %r414, %r415 }, { %r356, %r357 }; // 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 { %r360, %r361 }, { %r402, %r403, %r404, %r405 }, { %r416, %r417 }, { %r360, %r361 }; // 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 { %r364, %r365 }, { %r402, %r403, %r404, %r405 }, { %r418, %r419 }, { %r364, %r365 }; // 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 { %r368, %r369 }, { %r402, %r403, %r404, %r405 }, { %r420, %r421 }, { %r368, %r369 }; // 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 { %r372, %r373 }, { %r402, %r403, %r404, %r405 }, { %r422, %r423 }, { %r372, %r373 }; // 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 { %r376, %r377 }, { %r402, %r403, %r404, %r405 }, { %r424, %r425 }, { %r376, %r377 }; // 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 { %r380, %r381 }, { %r402, %r403, %r404, %r405 }, { %r426, %r427 }, { %r380, %r381 }; // 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 { %r384, %r385 }, { %r402, %r403, %r404, %r405 }, { %r428, %r429 }, { %r384, %r385 }; // 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 { %r388, %r389 }, { %r402, %r403, %r404, %r405 }, { %r430, %r431 }, { %r388, %r389 }; // 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 { %r392, %r393 }, { %r402, %r403, %r404, %r405 }, { %r432, %r433 }, { %r392, %r393 }; // 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 { %r396, %r397 }, { %r402, %r403, %r404, %r405 }, { %r434, %r435 }, { %r396, %r397 }; // 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 }, { %r438, %r439, %r440, %r441 }, { %r436, %r437 }, { %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 { %r336, %r337 }, { %r438, %r439, %r440, %r441 }, { %r442, %r443 }, { %r336, %r337 }; // 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 }, { %r438, %r439, %r440, %r441 }, { %r444, %r445 }, { %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 { %r348, %r349 }, { %r438, %r439, %r440, %r441 }, { %r446, %r447 }, { %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 { %r352, %r353 }, { %r438, %r439, %r440, %r441 }, { %r448, %r449 }, { %r352, %r353 }; // 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 { %r356, %r357 }, { %r438, %r439, %r440, %r441 }, { %r450, %r451 }, { %r356, %r357 }; // 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 { %r360, %r361 }, { %r438, %r439, %r440, %r441 }, { %r452, %r453 }, { %r360, %r361 }; // 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 { %r364, %r365 }, { %r438, %r439, %r440, %r441 }, { %r454, %r455 }, { %r364, %r365 }; // 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 { %r368, %r369 }, { %r438, %r439, %r440, %r441 }, { %r456, %r457 }, { %r368, %r369 }; // 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 { %r372, %r373 }, { %r438, %r439, %r440, %r441 }, { %r458, %r459 }, { %r372, %r373 }; // 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 { %r376, %r377 }, { %r438, %r439, %r440, %r441 }, { %r460, %r461 }, { %r376, %r377 }; // 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 { %r380, %r381 }, { %r438, %r439, %r440, %r441 }, { %r462, %r463 }, { %r380, %r381 }; // 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 { %r384, %r385 }, { %r438, %r439, %r440, %r441 }, { %r464, %r465 }, { %r384, %r385 }; // 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 { %r388, %r389 }, { %r438, %r439, %r440, %r441 }, { %r466, %r467 }, { %r388, %r389 }; // 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 { %r392, %r393 }, { %r438, %r439, %r440, %r441 }, { %r468, %r469 }, { %r392, %r393 }; // 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 { %r396, %r397 }, { %r438, %r439, %r440, %r441 }, { %r470, %r471 }, { %r396, %r397 }; // 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}, %r396; cvt.f32.f16 %r745, %rs1; cvt.f32.f16 %r746, %rs2; mov.b32 {%rs3, %rs4}, %r397; cvt.f32.f16 %r747, %rs3; cvt.f32.f16 %r748, %rs4; mov.b32 {%rs5, %rs6}, %r392; cvt.f32.f16 %r749, %rs5; cvt.f32.f16 %r750, %rs6; mov.b32 {%rs7, %rs8}, %r393; cvt.f32.f16 %r751, %rs7; cvt.f32.f16 %r752, %rs8; mov.b32 {%rs9, %rs10}, %r388; cvt.f32.f16 %r753, %rs9; cvt.f32.f16 %r754, %rs10; mov.b32 {%rs11, %rs12}, %r389; cvt.f32.f16 %r755, %rs11; cvt.f32.f16 %r756, %rs12; mov.b32 {%rs13, %rs14}, %r384; cvt.f32.f16 %r757, %rs13; cvt.f32.f16 %r758, %rs14; mov.b32 {%rs15, %rs16}, %r385; cvt.f32.f16 %r759, %rs15; cvt.f32.f16 %r760, %rs16; mov.b32 {%rs17, %rs18}, %r380; cvt.f32.f16 %r761, %rs17; cvt.f32.f16 %r762, %rs18; mov.b32 {%rs19, %rs20}, %r381; cvt.f32.f16 %r763, %rs19; cvt.f32.f16 %r764, %rs20; mov.b32 {%rs21, %rs22}, %r376; cvt.f32.f16 %r765, %rs21; cvt.f32.f16 %r766, %rs22; mov.b32 {%rs23, %rs24}, %r377; cvt.f32.f16 %r767, %rs23; cvt.f32.f16 %r768, %rs24; mov.b32 {%rs25, %rs26}, %r372; cvt.f32.f16 %r769, %rs25; cvt.f32.f16 %r770, %rs26; mov.b32 {%rs27, %rs28}, %r373; cvt.f32.f16 %r771, %rs27; cvt.f32.f16 %r772, %rs28; mov.b32 {%rs29, %rs30}, %r368; cvt.f32.f16 %r773, %rs29; cvt.f32.f16 %r774, %rs30; mov.b32 {%rs31, %rs32}, %r369; cvt.f32.f16 %r775, %rs31; cvt.f32.f16 %r776, %rs32; mov.b32 {%rs33, %rs34}, %r364; cvt.f32.f16 %r777, %rs33; cvt.f32.f16 %r778, %rs34; mov.b32 {%rs35, %rs36}, %r365; cvt.f32.f16 %r779, %rs35; cvt.f32.f16 %r780, %rs36; mov.b32 {%rs37, %rs38}, %r360; cvt.f32.f16 %r781, %rs37; cvt.f32.f16 %r782, %rs38; mov.b32 {%rs39, %rs40}, %r361; cvt.f32.f16 %r783, %rs39; cvt.f32.f16 %r784, %rs40; mov.b32 {%rs41, %rs42}, %r356; cvt.f32.f16 %r785, %rs41; cvt.f32.f16 %r786, %rs42; mov.b32 {%rs43, %rs44}, %r357; cvt.f32.f16 %r787, %rs43; cvt.f32.f16 %r788, %rs44; mov.b32 {%rs45, %rs46}, %r352; cvt.f32.f16 %r789, %rs45; cvt.f32.f16 %r790, %rs46; mov.b32 {%rs47, %rs48}, %r353; cvt.f32.f16 %r791, %rs47; cvt.f32.f16 %r792, %rs48; mov.b32 {%rs49, %rs50}, %r348; cvt.f32.f16 %r793, %rs49; cvt.f32.f16 %r794, %rs50; mov.b32 {%rs51, %rs52}, %r349; cvt.f32.f16 %r795, %rs51; cvt.f32.f16 %r796, %rs52; mov.b32 {%rs53, %rs54}, %r344; cvt.f32.f16 %r797, %rs53; cvt.f32.f16 %r798, %rs54; mov.b32 {%rs55, %rs56}, %r345; cvt.f32.f16 %r799, %rs55; cvt.f32.f16 %r800, %rs56; mov.b32 {%rs57, %rs58}, %r336; cvt.f32.f16 %r801, %rs57; cvt.f32.f16 %r802, %rs58; mov.b32 {%rs59, %rs60}, %r337; cvt.f32.f16 %r803, %rs59; cvt.f32.f16 %r804, %rs60; mov.b32 {%rs61, %rs62}, %r332; cvt.f32.f16 %r805, %rs61; cvt.f32.f16 %r806, %rs62; mov.b32 {%rs63, %rs64}, %r333; cvt.f32.f16 %r807, %rs63; cvt.f32.f16 %r808, %rs64; fma.rn.f32 %r1006, %r1006, %r735, %r808; fma.rn.f32 %r1005, %r1005, %r735, %r807; fma.rn.f32 %r1004, %r1004, %r734, %r806; fma.rn.f32 %r1003, %r1003, %r734, %r805; fma.rn.f32 %r1010, %r1010, %r735, %r804; fma.rn.f32 %r1009, %r1009, %r735, %r803; fma.rn.f32 %r1008, %r1008, %r734, %r802; fma.rn.f32 %r1007, %r1007, %r734, %r801; fma.rn.f32 %r1014, %r1014, %r735, %r800; fma.rn.f32 %r1013, %r1013, %r735, %r799; fma.rn.f32 %r1012, %r1012, %r734, %r798; fma.rn.f32 %r1011, %r1011, %r734, %r797; fma.rn.f32 %r1018, %r1018, %r735, %r796; fma.rn.f32 %r1017, %r1017, %r735, %r795; fma.rn.f32 %r1016, %r1016, %r734, %r794; fma.rn.f32 %r1015, %r1015, %r734, %r793; fma.rn.f32 %r1022, %r1022, %r735, %r792; fma.rn.f32 %r1021, %r1021, %r735, %r791; fma.rn.f32 %r1020, %r1020, %r734, %r790; fma.rn.f32 %r1019, %r1019, %r734, %r789; fma.rn.f32 %r1026, %r1026, %r735, %r788; fma.rn.f32 %r1025, %r1025, %r735, %r787; fma.rn.f32 %r1024, %r1024, %r734, %r786; fma.rn.f32 %r1023, %r1023, %r734, %r785; fma.rn.f32 %r1030, %r1030, %r735, %r784; fma.rn.f32 %r1029, %r1029, %r735, %r783; fma.rn.f32 %r1028, %r1028, %r734, %r782; fma.rn.f32 %r1027, %r1027, %r734, %r781; fma.rn.f32 %r1034, %r1034, %r735, %r780; fma.rn.f32 %r1033, %r1033, %r735, %r779; fma.rn.f32 %r1032, %r1032, %r734, %r778; fma.rn.f32 %r1031, %r1031, %r734, %r777; fma.rn.f32 %r1038, %r1038, %r735, %r776; fma.rn.f32 %r1037, %r1037, %r735, %r775; fma.rn.f32 %r1036, %r1036, %r734, %r774; fma.rn.f32 %r1035, %r1035, %r734, %r773; fma.rn.f32 %r1042, %r1042, %r735, %r772; fma.rn.f32 %r1041, %r1041, %r735, %r771; fma.rn.f32 %r1040, %r1040, %r734, %r770; fma.rn.f32 %r1039, %r1039, %r734, %r769; fma.rn.f32 %r1046, %r1046, %r735, %r768; fma.rn.f32 %r1045, %r1045, %r735, %r767; fma.rn.f32 %r1044, %r1044, %r734, %r766; fma.rn.f32 %r1043, %r1043, %r734, %r765; fma.rn.f32 %r1050, %r1050, %r735, %r764; fma.rn.f32 %r1049, %r1049, %r735, %r763; fma.rn.f32 %r1048, %r1048, %r734, %r762; fma.rn.f32 %r1047, %r1047, %r734, %r761; fma.rn.f32 %r1054, %r1054, %r735, %r760; fma.rn.f32 %r1053, %r1053, %r735, %r759; fma.rn.f32 %r1052, %r1052, %r734, %r758; fma.rn.f32 %r1051, %r1051, %r734, %r757; fma.rn.f32 %r1058, %r1058, %r735, %r756; fma.rn.f32 %r1057, %r1057, %r735, %r755; fma.rn.f32 %r1056, %r1056, %r734, %r754; fma.rn.f32 %r1055, %r1055, %r734, %r753; fma.rn.f32 %r1062, %r1062, %r735, %r752; fma.rn.f32 %r1061, %r1061, %r735, %r751; fma.rn.f32 %r1060, %r1060, %r734, %r750; fma.rn.f32 %r1059, %r1059, %r734, %r749; fma.rn.f32 %r1066, %r1066, %r735, %r748; fma.rn.f32 %r1065, %r1065, %r735, %r747; fma.rn.f32 %r1064, %r1064, %r734, %r746; fma.rn.f32 %r1063, %r1063, %r734, %r745; .loc 1 34 9, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:34:9 add.s64 %rd103, %rd135, %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 %rd104, %rd135, %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 %rd105, %rd136, %rd22; add.s64 %rd106, %rd136, %rd21; add.s64 %rd107, %rd136, %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 %rd108, %rd136, %rd18; add.s32 %r809, %r997, 1; setp.gt.s32 %p30, %r809, 2; selp.b32 %r997, 0, %r809, %p30; .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 %p31, %r3, %r995; setp.lt.s32 %p32, %r4, %r995; .loc 1 16 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:16:13 shl.b32 %r810, %r997, 13; add.s32 %r811, %r142, %r810; bar.sync 0; add.s32 %r812, %r13, %r810; add.s32 %r472, %r812, 49152; selp.b32 %r813, 16, 0, %p31; selp.b32 %r473, %r813, 0, %p28; // 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 [ %r472 + 0 ], [ %rd103 + 0 ], 0x10, %r473; // end inline asm add.s32 %r474, %r812, 53248; selp.b32 %r814, 16, 0, %p32; selp.b32 %r475, %r814, 0, %p28; // 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 [ %r474 + 0 ], [ %rd104 + 0 ], 0x10, %r475; // 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 %r815, %r997, 2; add.s32 %r816, %r142, %r815; add.s32 %r476, %r816, 90112; selp.b32 %r477, 4, 0, %p28; // begin inline asm @%p5 cp.async.ca.shared.global [ %r476 + 0 ], [ %rd134 + 0 ], 0x4, %r477; // 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 %p33, %r6, %r995; setp.lt.s32 %p34, %r7, %r995; setp.lt.s32 %p35, %r8, %r995; setp.lt.s32 %p36, %r9, %r995; .loc 1 29 13, function_name $L__info_string0, inlined_at 1 75 16 // attn_qk_int8_per_block.py:29:13 add.s32 %r817, %r811, %r810; add.s32 %r478, %r817, %r15; selp.b32 %r818, 16, 0, %p33; selp.b32 %r479, %r818, 0, %p28; // 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 [ %r478 + 0 ], [ %rd105 + 0 ], 0x10, %r479; // end inline asm add.s32 %r480, %r478, 4096; selp.b32 %r819, 16, 0, %p34; selp.b32 %r481, %r819, 0, %p28; // 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 [ %r480 + 0 ], [ %rd106 + 0 ], 0x10, %r481; // end inline asm add.s32 %r482, %r478, 8192; selp.b32 %r820, 16, 0, %p35; selp.b32 %r483, %r820, 0, %p28; // 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 [ %r482 + 0 ], [ %rd107 + 0 ], 0x10, %r483; // end inline asm add.s32 %r484, %r478, 12288; selp.b32 %r821, 16, 0, %p36; selp.b32 %r485, %r821, 0, %p28; // 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 [ %r484 + 0 ], [ %rd108 + 0 ], 0x10, %r485; // 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 %r1000, %r1000, 64; add.s64 %rd136, %rd136, %rd19; add.s64 %rd135, %rd135, %rd16; add.s64 %rd134, %rd134, 4; add.s32 %r995, %r995, -64; setp.lt.s32 %p37, %r1000, %r37; mov.b32 %r998, %r31; mov.b32 %r999, %r32; @%p37 bra $L__BB0_2; $L__tmp26: $L__BB0_3: // %._crit_edge .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r894, %r5, %r1; or.b32 %r895, %r894, 112; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p45, %r895, %r36; .loc 1 59 34 // attn_qk_int8_per_block.py:59:34 or.b32 %r896, %r6, %r1; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r897, %r896, 96; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p44, %r897, %r36; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r898, %r896, 80; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p43, %r898, %r36; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r899, %r896, 64; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p42, %r899, %r36; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r900, %r9, %r1; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p41, %r900, %r36; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r901, %r8, %r1; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p40, %r901, %r36; .loc 1 59 14 // attn_qk_int8_per_block.py:59:14 or.b32 %r902, %r7, %r1; .loc 1 73 32 // attn_qk_int8_per_block.py:73:32 setp.lt.s32 %p39, %r902, %r36; setp.lt.s32 %p38, %r896, %r36; .loc 1 67 26 // attn_qk_int8_per_block.py:67:26 cvt.s64.s32 %rd117, %r33; mul.lo.s64 %rd118, %rd117, %rd1; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 shl.b64 %rd119, %rd118, 1; add.s64 %rd120, %rd26, %rd119; .loc 1 67 46 // attn_qk_int8_per_block.py:67:46 cvt.s64.s32 %rd121, %r34; mul.lo.s64 %rd122, %rd121, %rd2; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 shl.b64 %rd123, %rd122, 1; add.s64 %rd124, %rd120, %rd123; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r903, %r35, %r895; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd125, %r903, 2, %rd124; shl.b64 %rd126, %rd15, 1; add.s64 %rd116, %rd125, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r904, %r35, %r897; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd127, %r904, 2, %rd124; add.s64 %rd115, %rd127, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r905, %r35, %r898; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd128, %r905, 2, %rd124; add.s64 %rd114, %rd128, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r906, %r35, %r899; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd129, %r906, 2, %rd124; add.s64 %rd113, %rd129, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r907, %r35, %r900; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd130, %r907, 2, %rd124; add.s64 %rd112, %rd130, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r908, %r35, %r901; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd131, %r908, 2, %rd124; add.s64 %rd111, %rd131, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r909, %r35, %r902; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd132, %r909, 2, %rd124; add.s64 %rd110, %rd132, %rd126; .loc 1 67 67 // attn_qk_int8_per_block.py:67:67 mul.lo.s32 %r910, %r35, %r896; .loc 1 67 19 // attn_qk_int8_per_block.py:67:19 mad.wide.s32 %rd133, %r910, 2, %rd124; add.s64 %rd109, %rd133, %rd126; $L__tmp27: .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__tmp28: .loc 1 80 11 // attn_qk_int8_per_block.py:80:11 div.full.f32 %r911, %r1003, %r1001; div.full.f32 %r912, %r1004, %r1001; div.full.f32 %r913, %r1005, %r1002; div.full.f32 %r914, %r1006, %r1002; div.full.f32 %r915, %r1007, %r1001; div.full.f32 %r916, %r1008, %r1001; div.full.f32 %r917, %r1009, %r1002; div.full.f32 %r918, %r1010, %r1002; div.full.f32 %r919, %r1011, %r1001; div.full.f32 %r920, %r1012, %r1001; div.full.f32 %r921, %r1013, %r1002; div.full.f32 %r922, %r1014, %r1002; div.full.f32 %r923, %r1015, %r1001; div.full.f32 %r924, %r1016, %r1001; div.full.f32 %r925, %r1017, %r1002; div.full.f32 %r926, %r1018, %r1002; div.full.f32 %r927, %r1019, %r1001; div.full.f32 %r928, %r1020, %r1001; div.full.f32 %r929, %r1021, %r1002; div.full.f32 %r930, %r1022, %r1002; div.full.f32 %r931, %r1023, %r1001; div.full.f32 %r932, %r1024, %r1001; div.full.f32 %r933, %r1025, %r1002; div.full.f32 %r934, %r1026, %r1002; div.full.f32 %r935, %r1027, %r1001; div.full.f32 %r936, %r1028, %r1001; div.full.f32 %r937, %r1029, %r1002; div.full.f32 %r938, %r1030, %r1002; div.full.f32 %r939, %r1031, %r1001; div.full.f32 %r940, %r1032, %r1001; div.full.f32 %r941, %r1033, %r1002; div.full.f32 %r942, %r1034, %r1002; div.full.f32 %r943, %r1035, %r1001; div.full.f32 %r944, %r1036, %r1001; div.full.f32 %r945, %r1037, %r1002; div.full.f32 %r946, %r1038, %r1002; div.full.f32 %r947, %r1039, %r1001; div.full.f32 %r948, %r1040, %r1001; div.full.f32 %r949, %r1041, %r1002; div.full.f32 %r950, %r1042, %r1002; div.full.f32 %r951, %r1043, %r1001; div.full.f32 %r952, %r1044, %r1001; div.full.f32 %r953, %r1045, %r1002; div.full.f32 %r954, %r1046, %r1002; div.full.f32 %r955, %r1047, %r1001; div.full.f32 %r956, %r1048, %r1001; div.full.f32 %r957, %r1049, %r1002; div.full.f32 %r958, %r1050, %r1002; div.full.f32 %r959, %r1051, %r1001; div.full.f32 %r960, %r1052, %r1001; div.full.f32 %r961, %r1053, %r1002; div.full.f32 %r962, %r1054, %r1002; div.full.f32 %r963, %r1055, %r1001; div.full.f32 %r964, %r1056, %r1001; div.full.f32 %r965, %r1057, %r1002; div.full.f32 %r966, %r1058, %r1002; div.full.f32 %r967, %r1059, %r1001; div.full.f32 %r968, %r1060, %r1001; div.full.f32 %r969, %r1061, %r1002; div.full.f32 %r970, %r1062, %r1002; div.full.f32 %r971, %r1063, %r1001; div.full.f32 %r972, %r1064, %r1001; div.full.f32 %r973, %r1065, %r1002; div.full.f32 %r974, %r1066, %r1002; .loc 1 81 27 // attn_qk_int8_per_block.py:81:27 cvt.rn.bf16x2.f32 %r823, %r912, %r911; cvt.rn.bf16x2.f32 %r828, %r914, %r913; cvt.rn.bf16x2.f32 %r824, %r916, %r915; cvt.rn.bf16x2.f32 %r829, %r918, %r917; cvt.rn.bf16x2.f32 %r825, %r920, %r919; cvt.rn.bf16x2.f32 %r830, %r922, %r921; cvt.rn.bf16x2.f32 %r826, %r924, %r923; cvt.rn.bf16x2.f32 %r831, %r926, %r925; cvt.rn.bf16x2.f32 %r833, %r928, %r927; cvt.rn.bf16x2.f32 %r838, %r930, %r929; cvt.rn.bf16x2.f32 %r834, %r932, %r931; cvt.rn.bf16x2.f32 %r839, %r934, %r933; cvt.rn.bf16x2.f32 %r835, %r936, %r935; cvt.rn.bf16x2.f32 %r840, %r938, %r937; cvt.rn.bf16x2.f32 %r836, %r940, %r939; cvt.rn.bf16x2.f32 %r841, %r942, %r941; cvt.rn.bf16x2.f32 %r843, %r944, %r943; cvt.rn.bf16x2.f32 %r848, %r946, %r945; cvt.rn.bf16x2.f32 %r844, %r948, %r947; cvt.rn.bf16x2.f32 %r849, %r950, %r949; cvt.rn.bf16x2.f32 %r845, %r952, %r951; cvt.rn.bf16x2.f32 %r850, %r954, %r953; cvt.rn.bf16x2.f32 %r846, %r956, %r955; cvt.rn.bf16x2.f32 %r851, %r958, %r957; cvt.rn.bf16x2.f32 %r853, %r960, %r959; cvt.rn.bf16x2.f32 %r858, %r962, %r961; cvt.rn.bf16x2.f32 %r854, %r964, %r963; cvt.rn.bf16x2.f32 %r859, %r966, %r965; cvt.rn.bf16x2.f32 %r855, %r968, %r967; cvt.rn.bf16x2.f32 %r860, %r970, %r969; cvt.rn.bf16x2.f32 %r856, %r972, %r971; cvt.rn.bf16x2.f32 %r861, %r974, %r973; .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 and.b32 %r975, %r2, 3; shl.b32 %r976, %r975, 13; shl.b32 %r977, %r2, 5; and.b32 %r978, %r977, 7264; and.b32 %r979, %r2, 24; shl.b32 %r980, %r979, 4; shl.b32 %r981, %r2, 2; and.b32 %r982, %r981, 16; or.b32 %r983, %r976, %r982; or.b32 %r984, %r978, %r980; or.b32 %r985, %r983, %r984; add.s32 %r822, %r142, %r985; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r822 + 0 ], { %r823, %r824, %r825, %r826 }; // end inline asm add.s32 %r827, %r822, 512; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r827 + 0 ], { %r828, %r829, %r830, %r831 }; // end inline asm xor.b32 %r986, %r985, 32; add.s32 %r832, %r142, %r986; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r832 + 0 ], { %r833, %r834, %r835, %r836 }; // end inline asm add.s32 %r837, %r832, 512; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r837 + 0 ], { %r838, %r839, %r840, %r841 }; // end inline asm xor.b32 %r987, %r985, 64; add.s32 %r842, %r142, %r987; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r842 + 0 ], { %r843, %r844, %r845, %r846 }; // end inline asm add.s32 %r847, %r842, 512; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r847 + 0 ], { %r848, %r849, %r850, %r851 }; // end inline asm xor.b32 %r988, %r985, 96; add.s32 %r852, %r142, %r988; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r852 + 0 ], { %r853, %r854, %r855, %r856 }; // end inline asm add.s32 %r857, %r852, 512; // begin inline asm .loc 1 81 5 // attn_qk_int8_per_block.py:81:5 st.shared::cta.v4.b32 [ %r857 + 0 ], { %r858, %r859, %r860, %r861 }; // end inline asm bar.sync 0; shl.b32 %r989, %r979, 10; shl.b32 %r990, %r975, 5; and.b32 %r991, %r981, 1008; or.b32 %r992, %r989, %r990; xor.b32 %r993, %r992, %r991; add.s32 %r994, %r142, %r993; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r862, %r863, %r864, %r865}, [%r994]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r866, %r867, %r868, %r869}, [%r994+1024]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r870, %r871, %r872, %r873}, [%r994+2048]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r874, %r875, %r876, %r877}, [%r994+3072]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r878, %r879, %r880, %r881}, [%r994+4096]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r882, %r883, %r884, %r885}, [%r994+5120]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r886, %r887, %r888, %r889}, [%r994+6144]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r890, %r891, %r892, %r893}, [%r994+7168]; // begin inline asm @%p38 st.global.v4.b32 [ %rd109 + 0 ], { %r862, %r863, %r864, %r865 }; // end inline asm // begin inline asm @%p39 st.global.v4.b32 [ %rd110 + 0 ], { %r866, %r867, %r868, %r869 }; // end inline asm // begin inline asm @%p40 st.global.v4.b32 [ %rd111 + 0 ], { %r870, %r871, %r872, %r873 }; // end inline asm // begin inline asm @%p41 st.global.v4.b32 [ %rd112 + 0 ], { %r874, %r875, %r876, %r877 }; // end inline asm // begin inline asm @%p42 st.global.v4.b32 [ %rd113 + 0 ], { %r878, %r879, %r880, %r881 }; // end inline asm // begin inline asm @%p43 st.global.v4.b32 [ %rd114 + 0 ], { %r882, %r883, %r884, %r885 }; // end inline asm // begin inline asm @%p44 st.global.v4.b32 [ %rd115 + 0 ], { %r886, %r887, %r888, %r889 }; // end inline asm // begin inline asm @%p45 st.global.v4.b32 [ %rd116 + 0 ], { %r890, %r891, %r892, %r893 }; // end inline asm .loc 1 40 1 // attn_qk_int8_per_block.py:40:1 ret; $L__tmp29: $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__tmp28 // 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__tmp15 // 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__tmp15 // 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__tmp16 // DW_AT_low_pc .b64 $L__tmp25 // 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__tmp16 // DW_AT_low_pc .b64 $L__tmp25 // 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 { }