Last active
June 5, 2026 22:08
-
-
Save malfet/510d239dc26a88643b4ec32fb0a7e840 to your computer and use it in GitHub Desktop.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // | |
| // Generated by LLVM NVPTX Back-End | |
| // | |
| .version 8.8 | |
| .target sm_100a | |
| .address_size 64 | |
| // .globl fused_moe_kernel // -- Begin function fused_moe_kernel | |
| .extern .shared .align 16 .b8 global_smem[]; | |
| // @fused_moe_kernel | |
| .visible .entry fused_moe_kernel( | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_0, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_1, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_2, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_3, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_4, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_5, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_6, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_7, | |
| .param .u32 fused_moe_kernel_param_8, | |
| .param .u32 fused_moe_kernel_param_9, | |
| .param .u32 fused_moe_kernel_param_10, | |
| .param .u32 fused_moe_kernel_param_11, | |
| .param .u32 fused_moe_kernel_param_12, | |
| .param .u32 fused_moe_kernel_param_13, | |
| .param .u32 fused_moe_kernel_param_14, | |
| .param .u32 fused_moe_kernel_param_15, | |
| .param .u32 fused_moe_kernel_param_16, | |
| .param .u32 fused_moe_kernel_param_17, | |
| .param .u32 fused_moe_kernel_param_18, | |
| .param .u32 fused_moe_kernel_param_19, | |
| .param .u32 fused_moe_kernel_param_20, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_21, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_22 | |
| ) | |
| .reqntid 128 | |
| { | |
| .reg .pred %p<82>; | |
| .reg .b32 %r<673>; | |
| .reg .b64 %rd<648>; | |
| .loc 1 47 0 // bar.py:47:0 | |
| $L__func_begin0: | |
| .loc 1 47 0 // bar.py:47:0 | |
| // %bb.0: | |
| ld.param.b32 %r19, [fused_moe_kernel_param_8]; | |
| $L__tmp0: | |
| .loc 1 47 0 // bar.py:47 | |
| mov.u32 %r1, %tid.x; | |
| setp.lt.u32 %p1, %r1, 32; | |
| mov.b32 %r29, global_smem; | |
| // begin inline asm | |
| @%p1 tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r29], 64; | |
| // end inline asm | |
| bar.sync 0; | |
| ld.shared.b32 %r2, [global_smem]; | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned; | |
| // end inline asm | |
| ld.param.b64 %rd120, [fused_moe_kernel_param_7]; | |
| .loc 1 93 24 // bar.py:93:24 | |
| mov.u32 %r31, %ctaid.x; | |
| ld.param.b32 %r32, [fused_moe_kernel_param_10]; | |
| $L__tmp1: | |
| .loc 2 43 17 // standard.py:43:17 @[ bar.py:94:28 ] | |
| add.s32 %r33, %r32, 63; | |
| .loc 2 43 30 // standard.py:43:30 @[ bar.py:94:28 ] | |
| shr.s32 %r34, %r33, 31; | |
| shr.u32 %r35, %r34, 26; | |
| add.s32 %r36, %r33, %r35; | |
| shr.s32 %r37, %r36, 6; | |
| $L__tmp2: | |
| .loc 2 43 17 // standard.py:43:17 @[ bar.py:95:27 ] | |
| add.s32 %r38, %r19, 63; | |
| .loc 2 43 30 // standard.py:43:30 @[ bar.py:95:27 ] | |
| shr.s32 %r39, %r38, 31; | |
| shr.u32 %r40, %r39, 26; | |
| add.s32 %r41, %r38, %r40; | |
| shr.s32 %r42, %r41, 6; | |
| $L__tmp3: | |
| .loc 1 96 38 // bar.py:96:38 | |
| shl.b32 %r43, %r42, 3; | |
| .loc 1 97 22 // bar.py:97:22 | |
| div.s32 %r44, %r31, %r43; | |
| .loc 1 98 29 // bar.py:98:29 | |
| shl.b32 %r45, %r44, 3; | |
| .loc 1 99 35 // bar.py:99:35 | |
| sub.s32 %r46, %r37, %r45; | |
| .loc 1 99 48 // bar.py:99:48 | |
| min.s32 %r47, %r46, 8; | |
| .loc 1 100 34 // bar.py:100:34 | |
| mul.lo.s32 %r48, %r44, %r43; | |
| sub.s32 %r49, %r31, %r48; | |
| .loc 1 101 40 // bar.py:101:40 | |
| div.s32 %r3, %r49, %r47; | |
| .loc 1 100 54 // bar.py:100:54 | |
| mul.lo.s32 %r50, %r3, %r47; | |
| sub.s32 %r51, %r49, %r50; | |
| .loc 1 100 27 // bar.py:100:27 | |
| add.s32 %r4, %r51, %r45; | |
| .loc 1 104 37 // bar.py:104:37 | |
| // begin inline asm | |
| mov.u32 %r30, 0x0; | |
| ld.global.b32 { %r30 }, [ %rd120 + 0 ]; | |
| // end inline asm | |
| .loc 1 105 15 // bar.py:105:15 | |
| shl.b32 %r10, %r4, 6; | |
| .loc 1 105 31 // bar.py:105:31 | |
| setp.lt.s32 %p2, %r10, %r30; | |
| @%p2 bra $L__BB0_3; | |
| bra.uni $L__BB0_1; | |
| $L__BB0_3: | |
| .loc 1 0 31 // bar.py:0:31 | |
| ld.param.b32 %r25, [fused_moe_kernel_param_15]; | |
| ld.param.b32 %r21, [fused_moe_kernel_param_11]; | |
| ld.param.b64 %rd119, [fused_moe_kernel_param_6]; | |
| ld.param.b64 %rd118, [fused_moe_kernel_param_5]; | |
| ld.param.b64 %rd115, [fused_moe_kernel_param_2]; | |
| .loc 1 103 0 // bar.py:103 | |
| bfe.u32 %r52, %r1, 1, 6; | |
| and.b32 %r6, %r1, 96; | |
| shl.b32 %r7, %r1, 3; | |
| and.b32 %r8, %r7, 56; | |
| cvt.u64.u32 %rd1, %r52; | |
| .loc 1 103 24 // bar.py:103:24 | |
| and.b32 %r11, %r1, 15; | |
| shr.u32 %r92, %r6, 1; | |
| or.b32 %r93, %r92, %r11; | |
| bfe.u32 %r94, %r1, 3, 4; | |
| .loc 1 109 52 // bar.py:109:52 | |
| mad.wide.s32 %rd128, %r10, 4, %rd118; | |
| mad.wide.u32 %rd121, %r94, 4, %rd128; | |
| add.s64 %rd122, %rd121, 64; | |
| add.s64 %rd123, %rd121, 128; | |
| add.s64 %rd124, %rd121, 192; | |
| shl.b64 %rd129, %rd1, 2; | |
| add.s64 %rd125, %rd128, %rd129; | |
| mad.wide.u32 %rd126, %r93, 4, %rd128; | |
| .loc 1 109 29 // bar.py:109:29 | |
| // begin inline asm | |
| mov.u32 %r85, 0x0; | |
| ld.global.b32 { %r85 }, [ %rd121 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r86, 0x0; | |
| ld.global.b32 { %r86 }, [ %rd122 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r87, 0x0; | |
| ld.global.b32 { %r87 }, [ %rd123 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r88, 0x0; | |
| ld.global.b32 { %r88 }, [ %rd124 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r89, 0x0; | |
| ld.global.b32 { %r89 }, [ %rd125 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r90, 0x0; | |
| ld.global.b32 { %r90 }, [ %rd126 + 0 ]; | |
| // end inline asm | |
| .loc 1 112 31 // bar.py:112:31 | |
| cvt.s64.s32 %rd34, %r85; | |
| cvt.s64.s32 %rd35, %r86; | |
| cvt.s64.s32 %rd36, %r87; | |
| cvt.s64.s32 %rd37, %r88; | |
| .loc 1 116 43 // bar.py:116:43 | |
| mad.wide.s32 %rd127, %r4, 4, %rd119; | |
| .loc 1 116 26 // bar.py:116:26 | |
| // begin inline asm | |
| mov.u32 %r91, 0x0; | |
| ld.global.b32 { %r91 }, [ %rd127 + 0 ]; | |
| // end inline asm | |
| .loc 1 117 22 // bar.py:117:22 | |
| setp.ne.b32 %p3, %r91, -1; | |
| shl.b32 %r663, %r3, 6; | |
| @%p3 bra $L__BB0_5; | |
| bra.uni $L__BB0_4; | |
| $L__BB0_5: | |
| .loc 1 0 0 // bar.py:0 | |
| and.b32 %r9, %r1, 16; | |
| shl.b32 %r53, %r9, 1; | |
| or.b32 %r54, %r53, 1; | |
| or.b32 %r55, %r53, 2; | |
| cvt.u64.u32 %rd2, %r53; | |
| cvt.u64.u32 %rd3, %r54; | |
| .loc 1 132 38 // bar.py:132:38 | |
| cvt.s64.s32 %rd130, %r663; | |
| or.b64 %rd39, %rd130, %rd1; | |
| or.b64 %rd40, %rd130, %rd2; | |
| .loc 1 132 81 // bar.py:132:81 | |
| cvt.s64.s32 %rd72, %r19; | |
| or.b64 %rd131, %rd130, %rd72; | |
| and.b64 %rd132, %rd131, -4294967296; | |
| setp.ne.b64 %p4, %rd132, 0; | |
| @%p4 bra $L__BB0_7; | |
| bra.uni $L__BB0_6; | |
| $L__BB0_7: | |
| rem.s64 %rd597, %rd39, %rd72; | |
| bra.uni $L__BB0_8; | |
| $L__BB0_1: | |
| .loc 1 106 8 // bar.py:106:8 | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r2, 64; | |
| // end inline asm | |
| bra.uni $L__BB0_2; | |
| $L__BB0_4: | |
| .loc 1 0 8 // bar.py:0:8 | |
| cvt.u32.u64 %r656, %rd37; | |
| cvt.u32.u64 %r657, %rd36; | |
| cvt.u32.u64 %r658, %rd35; | |
| cvt.u32.u64 %r659, %rd34; | |
| .loc 1 114 30 // bar.py:114:30 | |
| setp.lt.s32 %p74, %r656, %r21; | |
| setp.lt.s32 %p75, %r657, %r21; | |
| setp.lt.s32 %p76, %r658, %r21; | |
| setp.lt.s32 %p77, %r659, %r21; | |
| $L__tmp4: | |
| .loc 1 40 37 // bar.py:40:37 @[ bar.py:128:12 ] | |
| or.b32 %r660, %r663, %r8; | |
| .loc 1 41 33 // bar.py:41:33 @[ bar.py:128:12 ] | |
| cvt.s64.s32 %rd583, %r25; | |
| mul.lo.s64 %rd584, %rd34, %rd583; | |
| mul.lo.s64 %rd585, %rd35, %rd583; | |
| mul.lo.s64 %rd586, %rd36, %rd583; | |
| mul.lo.s64 %rd587, %rd37, %rd583; | |
| .loc 1 41 21 // bar.py:41:21 @[ bar.py:128:12 ] | |
| shl.b64 %rd588, %rd584, 1; | |
| add.s64 %rd589, %rd115, %rd588; | |
| shl.b64 %rd590, %rd585, 1; | |
| add.s64 %rd591, %rd115, %rd590; | |
| shl.b64 %rd592, %rd586, 1; | |
| add.s64 %rd593, %rd115, %rd592; | |
| shl.b64 %rd594, %rd587, 1; | |
| add.s64 %rd595, %rd115, %rd594; | |
| .loc 1 41 55 // bar.py:41:55 @[ bar.py:128:12 ] | |
| mul.wide.s32 %rd596, %r660, 2; | |
| add.s64 %rd579, %rd589, %rd596; | |
| add.s64 %rd580, %rd591, %rd596; | |
| add.s64 %rd581, %rd593, %rd596; | |
| add.s64 %rd582, %rd595, %rd596; | |
| .loc 1 42 55 // bar.py:42:55 @[ bar.py:128:12 ] | |
| setp.lt.s32 %p78, %r660, %r19; | |
| .loc 1 42 36 // bar.py:42:36 @[ bar.py:128:12 ] | |
| and.pred %p70, %p78, %p77; | |
| and.pred %p71, %p78, %p76; | |
| and.pred %p72, %p78, %p75; | |
| and.pred %p73, %p78, %p74; | |
| mov.b32 %r655, 0; | |
| .loc 1 43 21 // bar.py:43:21 @[ bar.py:128:12 ] | |
| // begin inline asm | |
| @%p70 st.global.v4.b32 [ %rd579 + 0 ], { %r655, %r655, %r655, %r655 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p71 st.global.v4.b32 [ %rd580 + 0 ], { %r655, %r655, %r655, %r655 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p72 st.global.v4.b32 [ %rd581 + 0 ], { %r655, %r655, %r655, %r655 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p73 st.global.v4.b32 [ %rd582 + 0 ], { %r655, %r655, %r655, %r655 }; | |
| // end inline asm | |
| $L__tmp5: | |
| .loc 1 130 8 // bar.py:130:8 | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r2, 64; | |
| // end inline asm | |
| bra.uni $L__BB0_2; | |
| $L__BB0_6: | |
| .loc 1 132 81 // bar.py:132:81 | |
| cvt.u32.u64 %r95, %rd72; | |
| cvt.u32.u64 %r96, %rd39; | |
| rem.u32 %r97, %r96, %r95; | |
| cvt.u64.u32 %rd597, %r97; | |
| $L__BB0_8: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r56, %r53, 3; | |
| cvt.u64.u32 %rd4, %r55; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd41, %rd130, %rd3; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd133, %rd40, %rd72; | |
| and.b64 %rd134, %rd133, -4294967296; | |
| setp.ne.b64 %p5, %rd134, 0; | |
| @%p5 bra $L__BB0_10; | |
| bra.uni $L__BB0_9; | |
| $L__BB0_10: | |
| rem.s64 %rd598, %rd40, %rd72; | |
| bra.uni $L__BB0_11; | |
| $L__BB0_9: | |
| cvt.u32.u64 %r98, %rd72; | |
| cvt.u32.u64 %r99, %rd40; | |
| rem.u32 %r100, %r99, %r98; | |
| cvt.u64.u32 %rd598, %r100; | |
| $L__BB0_11: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r57, %r53, 4; | |
| cvt.u64.u32 %rd5, %r56; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd42, %rd130, %rd4; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd135, %rd41, %rd72; | |
| and.b64 %rd136, %rd135, -4294967296; | |
| setp.ne.b64 %p6, %rd136, 0; | |
| @%p6 bra $L__BB0_13; | |
| bra.uni $L__BB0_12; | |
| $L__BB0_13: | |
| rem.s64 %rd599, %rd41, %rd72; | |
| bra.uni $L__BB0_14; | |
| $L__BB0_12: | |
| cvt.u32.u64 %r101, %rd72; | |
| cvt.u32.u64 %r102, %rd41; | |
| rem.u32 %r103, %r102, %r101; | |
| cvt.u64.u32 %rd599, %r103; | |
| $L__BB0_14: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r58, %r53, 5; | |
| cvt.u64.u32 %rd6, %r57; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd43, %rd130, %rd5; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd137, %rd42, %rd72; | |
| and.b64 %rd138, %rd137, -4294967296; | |
| setp.ne.b64 %p7, %rd138, 0; | |
| @%p7 bra $L__BB0_16; | |
| bra.uni $L__BB0_15; | |
| $L__BB0_16: | |
| rem.s64 %rd600, %rd42, %rd72; | |
| bra.uni $L__BB0_17; | |
| $L__BB0_15: | |
| cvt.u32.u64 %r104, %rd72; | |
| cvt.u32.u64 %r105, %rd42; | |
| rem.u32 %r106, %r105, %r104; | |
| cvt.u64.u32 %rd600, %r106; | |
| $L__BB0_17: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r59, %r53, 6; | |
| cvt.u64.u32 %rd7, %r58; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd44, %rd130, %rd6; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd139, %rd43, %rd72; | |
| and.b64 %rd140, %rd139, -4294967296; | |
| setp.ne.b64 %p8, %rd140, 0; | |
| @%p8 bra $L__BB0_19; | |
| bra.uni $L__BB0_18; | |
| $L__BB0_19: | |
| rem.s64 %rd601, %rd43, %rd72; | |
| bra.uni $L__BB0_20; | |
| $L__BB0_18: | |
| cvt.u32.u64 %r107, %rd72; | |
| cvt.u32.u64 %r108, %rd43; | |
| rem.u32 %r109, %r108, %r107; | |
| cvt.u64.u32 %rd601, %r109; | |
| $L__BB0_20: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r60, %r53, 7; | |
| cvt.u64.u32 %rd8, %r59; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd45, %rd130, %rd7; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd141, %rd44, %rd72; | |
| and.b64 %rd142, %rd141, -4294967296; | |
| setp.ne.b64 %p9, %rd142, 0; | |
| @%p9 bra $L__BB0_22; | |
| bra.uni $L__BB0_21; | |
| $L__BB0_22: | |
| rem.s64 %rd602, %rd44, %rd72; | |
| bra.uni $L__BB0_23; | |
| $L__BB0_21: | |
| cvt.u32.u64 %r110, %rd72; | |
| cvt.u32.u64 %r111, %rd44; | |
| rem.u32 %r112, %r111, %r110; | |
| cvt.u64.u32 %rd602, %r112; | |
| $L__BB0_23: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r61, %r53, 8; | |
| cvt.u64.u32 %rd9, %r60; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd46, %rd130, %rd8; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd143, %rd45, %rd72; | |
| and.b64 %rd144, %rd143, -4294967296; | |
| setp.ne.b64 %p10, %rd144, 0; | |
| @%p10 bra $L__BB0_25; | |
| bra.uni $L__BB0_24; | |
| $L__BB0_25: | |
| rem.s64 %rd603, %rd45, %rd72; | |
| bra.uni $L__BB0_26; | |
| $L__BB0_24: | |
| cvt.u32.u64 %r113, %rd72; | |
| cvt.u32.u64 %r114, %rd45; | |
| rem.u32 %r115, %r114, %r113; | |
| cvt.u64.u32 %rd603, %r115; | |
| $L__BB0_26: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r62, %r53, 9; | |
| cvt.u64.u32 %rd10, %r61; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd47, %rd130, %rd9; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd145, %rd46, %rd72; | |
| and.b64 %rd146, %rd145, -4294967296; | |
| setp.ne.b64 %p11, %rd146, 0; | |
| @%p11 bra $L__BB0_28; | |
| bra.uni $L__BB0_27; | |
| $L__BB0_28: | |
| rem.s64 %rd604, %rd46, %rd72; | |
| bra.uni $L__BB0_29; | |
| $L__BB0_27: | |
| cvt.u32.u64 %r116, %rd72; | |
| cvt.u32.u64 %r117, %rd46; | |
| rem.u32 %r118, %r117, %r116; | |
| cvt.u64.u32 %rd604, %r118; | |
| $L__BB0_29: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r63, %r53, 10; | |
| cvt.u64.u32 %rd11, %r62; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd48, %rd130, %rd10; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd147, %rd47, %rd72; | |
| and.b64 %rd148, %rd147, -4294967296; | |
| setp.ne.b64 %p12, %rd148, 0; | |
| @%p12 bra $L__BB0_31; | |
| bra.uni $L__BB0_30; | |
| $L__BB0_31: | |
| rem.s64 %rd605, %rd47, %rd72; | |
| bra.uni $L__BB0_32; | |
| $L__BB0_30: | |
| cvt.u32.u64 %r119, %rd72; | |
| cvt.u32.u64 %r120, %rd47; | |
| rem.u32 %r121, %r120, %r119; | |
| cvt.u64.u32 %rd605, %r121; | |
| $L__BB0_32: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r64, %r53, 11; | |
| cvt.u64.u32 %rd12, %r63; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd49, %rd130, %rd11; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd149, %rd48, %rd72; | |
| and.b64 %rd150, %rd149, -4294967296; | |
| setp.ne.b64 %p13, %rd150, 0; | |
| @%p13 bra $L__BB0_34; | |
| bra.uni $L__BB0_33; | |
| $L__BB0_34: | |
| rem.s64 %rd606, %rd48, %rd72; | |
| bra.uni $L__BB0_35; | |
| $L__BB0_33: | |
| cvt.u32.u64 %r122, %rd72; | |
| cvt.u32.u64 %r123, %rd48; | |
| rem.u32 %r124, %r123, %r122; | |
| cvt.u64.u32 %rd606, %r124; | |
| $L__BB0_35: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r65, %r53, 12; | |
| cvt.u64.u32 %rd13, %r64; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd50, %rd130, %rd12; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd151, %rd49, %rd72; | |
| and.b64 %rd152, %rd151, -4294967296; | |
| setp.ne.b64 %p14, %rd152, 0; | |
| @%p14 bra $L__BB0_37; | |
| bra.uni $L__BB0_36; | |
| $L__BB0_37: | |
| rem.s64 %rd607, %rd49, %rd72; | |
| bra.uni $L__BB0_38; | |
| $L__BB0_36: | |
| cvt.u32.u64 %r125, %rd72; | |
| cvt.u32.u64 %r126, %rd49; | |
| rem.u32 %r127, %r126, %r125; | |
| cvt.u64.u32 %rd607, %r127; | |
| $L__BB0_38: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r66, %r53, 13; | |
| cvt.u64.u32 %rd14, %r65; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd51, %rd130, %rd13; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd153, %rd50, %rd72; | |
| and.b64 %rd154, %rd153, -4294967296; | |
| setp.ne.b64 %p15, %rd154, 0; | |
| @%p15 bra $L__BB0_40; | |
| bra.uni $L__BB0_39; | |
| $L__BB0_40: | |
| rem.s64 %rd608, %rd50, %rd72; | |
| bra.uni $L__BB0_41; | |
| $L__BB0_39: | |
| cvt.u32.u64 %r128, %rd72; | |
| cvt.u32.u64 %r129, %rd50; | |
| rem.u32 %r130, %r129, %r128; | |
| cvt.u64.u32 %rd608, %r130; | |
| $L__BB0_41: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r67, %r53, 14; | |
| cvt.u64.u32 %rd15, %r66; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd52, %rd130, %rd14; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd155, %rd51, %rd72; | |
| and.b64 %rd156, %rd155, -4294967296; | |
| setp.ne.b64 %p16, %rd156, 0; | |
| @%p16 bra $L__BB0_43; | |
| bra.uni $L__BB0_42; | |
| $L__BB0_43: | |
| rem.s64 %rd609, %rd51, %rd72; | |
| bra.uni $L__BB0_44; | |
| $L__BB0_42: | |
| cvt.u32.u64 %r131, %rd72; | |
| cvt.u32.u64 %r132, %rd51; | |
| rem.u32 %r133, %r132, %r131; | |
| cvt.u64.u32 %rd609, %r133; | |
| $L__BB0_44: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r68, %r53, 15; | |
| cvt.u64.u32 %rd16, %r67; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd53, %rd130, %rd15; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd157, %rd52, %rd72; | |
| and.b64 %rd158, %rd157, -4294967296; | |
| setp.ne.b64 %p17, %rd158, 0; | |
| @%p17 bra $L__BB0_46; | |
| bra.uni $L__BB0_45; | |
| $L__BB0_46: | |
| rem.s64 %rd610, %rd52, %rd72; | |
| bra.uni $L__BB0_47; | |
| $L__BB0_45: | |
| cvt.u32.u64 %r134, %rd72; | |
| cvt.u32.u64 %r135, %rd52; | |
| rem.u32 %r136, %r135, %r134; | |
| cvt.u64.u32 %rd610, %r136; | |
| $L__BB0_47: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r69, %r53, 16; | |
| cvt.u64.u32 %rd17, %r68; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd54, %rd130, %rd16; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd159, %rd53, %rd72; | |
| and.b64 %rd160, %rd159, -4294967296; | |
| setp.ne.b64 %p18, %rd160, 0; | |
| @%p18 bra $L__BB0_49; | |
| bra.uni $L__BB0_48; | |
| $L__BB0_49: | |
| rem.s64 %rd611, %rd53, %rd72; | |
| bra.uni $L__BB0_50; | |
| $L__BB0_48: | |
| cvt.u32.u64 %r137, %rd72; | |
| cvt.u32.u64 %r138, %rd53; | |
| rem.u32 %r139, %r138, %r137; | |
| cvt.u64.u32 %rd611, %r139; | |
| $L__BB0_50: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r70, %r53, 17; | |
| cvt.u64.u32 %rd18, %r69; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd55, %rd130, %rd17; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd161, %rd54, %rd72; | |
| and.b64 %rd162, %rd161, -4294967296; | |
| setp.ne.b64 %p19, %rd162, 0; | |
| @%p19 bra $L__BB0_52; | |
| bra.uni $L__BB0_51; | |
| $L__BB0_52: | |
| rem.s64 %rd612, %rd54, %rd72; | |
| bra.uni $L__BB0_53; | |
| $L__BB0_51: | |
| cvt.u32.u64 %r140, %rd72; | |
| cvt.u32.u64 %r141, %rd54; | |
| rem.u32 %r142, %r141, %r140; | |
| cvt.u64.u32 %rd612, %r142; | |
| $L__BB0_53: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r71, %r53, 18; | |
| cvt.u64.u32 %rd19, %r70; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd56, %rd130, %rd18; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd163, %rd55, %rd72; | |
| and.b64 %rd164, %rd163, -4294967296; | |
| setp.ne.b64 %p20, %rd164, 0; | |
| @%p20 bra $L__BB0_55; | |
| bra.uni $L__BB0_54; | |
| $L__BB0_55: | |
| rem.s64 %rd613, %rd55, %rd72; | |
| bra.uni $L__BB0_56; | |
| $L__BB0_54: | |
| cvt.u32.u64 %r143, %rd72; | |
| cvt.u32.u64 %r144, %rd55; | |
| rem.u32 %r145, %r144, %r143; | |
| cvt.u64.u32 %rd613, %r145; | |
| $L__BB0_56: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r72, %r53, 19; | |
| cvt.u64.u32 %rd20, %r71; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd57, %rd130, %rd19; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd165, %rd56, %rd72; | |
| and.b64 %rd166, %rd165, -4294967296; | |
| setp.ne.b64 %p21, %rd166, 0; | |
| @%p21 bra $L__BB0_58; | |
| bra.uni $L__BB0_57; | |
| $L__BB0_58: | |
| rem.s64 %rd614, %rd56, %rd72; | |
| bra.uni $L__BB0_59; | |
| $L__BB0_57: | |
| cvt.u32.u64 %r146, %rd72; | |
| cvt.u32.u64 %r147, %rd56; | |
| rem.u32 %r148, %r147, %r146; | |
| cvt.u64.u32 %rd614, %r148; | |
| $L__BB0_59: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r73, %r53, 20; | |
| cvt.u64.u32 %rd21, %r72; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd58, %rd130, %rd20; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd167, %rd57, %rd72; | |
| and.b64 %rd168, %rd167, -4294967296; | |
| setp.ne.b64 %p22, %rd168, 0; | |
| @%p22 bra $L__BB0_61; | |
| bra.uni $L__BB0_60; | |
| $L__BB0_61: | |
| rem.s64 %rd615, %rd57, %rd72; | |
| bra.uni $L__BB0_62; | |
| $L__BB0_60: | |
| cvt.u32.u64 %r149, %rd72; | |
| cvt.u32.u64 %r150, %rd57; | |
| rem.u32 %r151, %r150, %r149; | |
| cvt.u64.u32 %rd615, %r151; | |
| $L__BB0_62: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r74, %r53, 21; | |
| cvt.u64.u32 %rd22, %r73; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd59, %rd130, %rd21; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd169, %rd58, %rd72; | |
| and.b64 %rd170, %rd169, -4294967296; | |
| setp.ne.b64 %p23, %rd170, 0; | |
| @%p23 bra $L__BB0_64; | |
| bra.uni $L__BB0_63; | |
| $L__BB0_64: | |
| rem.s64 %rd616, %rd58, %rd72; | |
| bra.uni $L__BB0_65; | |
| $L__BB0_63: | |
| cvt.u32.u64 %r152, %rd72; | |
| cvt.u32.u64 %r153, %rd58; | |
| rem.u32 %r154, %r153, %r152; | |
| cvt.u64.u32 %rd616, %r154; | |
| $L__BB0_65: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r75, %r53, 22; | |
| cvt.u64.u32 %rd23, %r74; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd60, %rd130, %rd22; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd171, %rd59, %rd72; | |
| and.b64 %rd172, %rd171, -4294967296; | |
| setp.ne.b64 %p24, %rd172, 0; | |
| @%p24 bra $L__BB0_67; | |
| bra.uni $L__BB0_66; | |
| $L__BB0_67: | |
| rem.s64 %rd617, %rd59, %rd72; | |
| bra.uni $L__BB0_68; | |
| $L__BB0_66: | |
| cvt.u32.u64 %r155, %rd72; | |
| cvt.u32.u64 %r156, %rd59; | |
| rem.u32 %r157, %r156, %r155; | |
| cvt.u64.u32 %rd617, %r157; | |
| $L__BB0_68: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r76, %r53, 23; | |
| cvt.u64.u32 %rd24, %r75; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd61, %rd130, %rd23; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd173, %rd60, %rd72; | |
| and.b64 %rd174, %rd173, -4294967296; | |
| setp.ne.b64 %p25, %rd174, 0; | |
| @%p25 bra $L__BB0_70; | |
| bra.uni $L__BB0_69; | |
| $L__BB0_70: | |
| rem.s64 %rd618, %rd60, %rd72; | |
| bra.uni $L__BB0_71; | |
| $L__BB0_69: | |
| cvt.u32.u64 %r158, %rd72; | |
| cvt.u32.u64 %r159, %rd60; | |
| rem.u32 %r160, %r159, %r158; | |
| cvt.u64.u32 %rd618, %r160; | |
| $L__BB0_71: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r77, %r53, 24; | |
| cvt.u64.u32 %rd25, %r76; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd62, %rd130, %rd24; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd175, %rd61, %rd72; | |
| and.b64 %rd176, %rd175, -4294967296; | |
| setp.ne.b64 %p26, %rd176, 0; | |
| @%p26 bra $L__BB0_73; | |
| bra.uni $L__BB0_72; | |
| $L__BB0_73: | |
| rem.s64 %rd619, %rd61, %rd72; | |
| bra.uni $L__BB0_74; | |
| $L__BB0_72: | |
| cvt.u32.u64 %r161, %rd72; | |
| cvt.u32.u64 %r162, %rd61; | |
| rem.u32 %r163, %r162, %r161; | |
| cvt.u64.u32 %rd619, %r163; | |
| $L__BB0_74: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r78, %r53, 25; | |
| cvt.u64.u32 %rd26, %r77; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd63, %rd130, %rd25; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd177, %rd62, %rd72; | |
| and.b64 %rd178, %rd177, -4294967296; | |
| setp.ne.b64 %p27, %rd178, 0; | |
| @%p27 bra $L__BB0_76; | |
| bra.uni $L__BB0_75; | |
| $L__BB0_76: | |
| rem.s64 %rd620, %rd62, %rd72; | |
| bra.uni $L__BB0_77; | |
| $L__BB0_75: | |
| cvt.u32.u64 %r164, %rd72; | |
| cvt.u32.u64 %r165, %rd62; | |
| rem.u32 %r166, %r165, %r164; | |
| cvt.u64.u32 %rd620, %r166; | |
| $L__BB0_77: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r79, %r53, 26; | |
| cvt.u64.u32 %rd27, %r78; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd64, %rd130, %rd26; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd179, %rd63, %rd72; | |
| and.b64 %rd180, %rd179, -4294967296; | |
| setp.ne.b64 %p28, %rd180, 0; | |
| @%p28 bra $L__BB0_79; | |
| bra.uni $L__BB0_78; | |
| $L__BB0_79: | |
| rem.s64 %rd621, %rd63, %rd72; | |
| bra.uni $L__BB0_80; | |
| $L__BB0_78: | |
| cvt.u32.u64 %r167, %rd72; | |
| cvt.u32.u64 %r168, %rd63; | |
| rem.u32 %r169, %r168, %r167; | |
| cvt.u64.u32 %rd621, %r169; | |
| $L__BB0_80: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r80, %r53, 27; | |
| cvt.u64.u32 %rd28, %r79; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd65, %rd130, %rd27; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd181, %rd64, %rd72; | |
| and.b64 %rd182, %rd181, -4294967296; | |
| setp.ne.b64 %p29, %rd182, 0; | |
| @%p29 bra $L__BB0_82; | |
| bra.uni $L__BB0_81; | |
| $L__BB0_82: | |
| rem.s64 %rd622, %rd64, %rd72; | |
| bra.uni $L__BB0_83; | |
| $L__BB0_81: | |
| cvt.u32.u64 %r170, %rd72; | |
| cvt.u32.u64 %r171, %rd64; | |
| rem.u32 %r172, %r171, %r170; | |
| cvt.u64.u32 %rd622, %r172; | |
| $L__BB0_83: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r81, %r53, 28; | |
| cvt.u64.u32 %rd29, %r80; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd66, %rd130, %rd28; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd183, %rd65, %rd72; | |
| and.b64 %rd184, %rd183, -4294967296; | |
| setp.ne.b64 %p30, %rd184, 0; | |
| @%p30 bra $L__BB0_85; | |
| bra.uni $L__BB0_84; | |
| $L__BB0_85: | |
| rem.s64 %rd623, %rd65, %rd72; | |
| bra.uni $L__BB0_86; | |
| $L__BB0_84: | |
| cvt.u32.u64 %r173, %rd72; | |
| cvt.u32.u64 %r174, %rd65; | |
| rem.u32 %r175, %r174, %r173; | |
| cvt.u64.u32 %rd623, %r175; | |
| $L__BB0_86: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r82, %r53, 29; | |
| cvt.u64.u32 %rd30, %r81; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd67, %rd130, %rd29; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd185, %rd66, %rd72; | |
| and.b64 %rd186, %rd185, -4294967296; | |
| setp.ne.b64 %p31, %rd186, 0; | |
| @%p31 bra $L__BB0_88; | |
| bra.uni $L__BB0_87; | |
| $L__BB0_88: | |
| rem.s64 %rd624, %rd66, %rd72; | |
| bra.uni $L__BB0_89; | |
| $L__BB0_87: | |
| cvt.u32.u64 %r176, %rd72; | |
| cvt.u32.u64 %r177, %rd66; | |
| rem.u32 %r178, %r177, %r176; | |
| cvt.u64.u32 %rd624, %r178; | |
| $L__BB0_89: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r83, %r53, 30; | |
| cvt.u64.u32 %rd31, %r82; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd68, %rd130, %rd30; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd187, %rd67, %rd72; | |
| and.b64 %rd188, %rd187, -4294967296; | |
| setp.ne.b64 %p32, %rd188, 0; | |
| @%p32 bra $L__BB0_91; | |
| bra.uni $L__BB0_90; | |
| $L__BB0_91: | |
| rem.s64 %rd625, %rd67, %rd72; | |
| bra.uni $L__BB0_92; | |
| $L__BB0_90: | |
| cvt.u32.u64 %r179, %rd72; | |
| cvt.u32.u64 %r180, %rd67; | |
| rem.u32 %r181, %r180, %r179; | |
| cvt.u64.u32 %rd625, %r181; | |
| $L__BB0_92: | |
| .loc 1 0 0 // bar.py:0 | |
| or.b32 %r84, %r53, 31; | |
| cvt.u64.u32 %rd32, %r83; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd69, %rd130, %rd31; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd189, %rd68, %rd72; | |
| and.b64 %rd190, %rd189, -4294967296; | |
| setp.ne.b64 %p33, %rd190, 0; | |
| @%p33 bra $L__BB0_94; | |
| bra.uni $L__BB0_93; | |
| $L__BB0_94: | |
| rem.s64 %rd626, %rd68, %rd72; | |
| bra.uni $L__BB0_95; | |
| $L__BB0_93: | |
| cvt.u32.u64 %r182, %rd72; | |
| cvt.u32.u64 %r183, %rd68; | |
| rem.u32 %r184, %r183, %r182; | |
| cvt.u64.u32 %rd626, %r184; | |
| $L__BB0_95: | |
| .loc 1 0 0 // bar.py:0 | |
| cvt.u64.u32 %rd33, %r84; | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd70, %rd130, %rd32; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd191, %rd69, %rd72; | |
| and.b64 %rd192, %rd191, -4294967296; | |
| setp.ne.b64 %p34, %rd192, 0; | |
| @%p34 bra $L__BB0_97; | |
| bra.uni $L__BB0_96; | |
| $L__BB0_97: | |
| rem.s64 %rd627, %rd69, %rd72; | |
| bra.uni $L__BB0_98; | |
| $L__BB0_96: | |
| cvt.u32.u64 %r185, %rd72; | |
| cvt.u32.u64 %r186, %rd69; | |
| rem.u32 %r187, %r186, %r185; | |
| cvt.u64.u32 %rd627, %r187; | |
| $L__BB0_98: | |
| .loc 1 132 0 // bar.py:132 | |
| or.b64 %rd71, %rd130, %rd33; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd193, %rd70, %rd72; | |
| and.b64 %rd194, %rd193, -4294967296; | |
| setp.ne.b64 %p35, %rd194, 0; | |
| @%p35 bra $L__BB0_100; | |
| bra.uni $L__BB0_99; | |
| $L__BB0_100: | |
| rem.s64 %rd628, %rd70, %rd72; | |
| bra.uni $L__BB0_101; | |
| $L__BB0_99: | |
| cvt.u32.u64 %r188, %rd72; | |
| cvt.u32.u64 %r189, %rd70; | |
| rem.u32 %r190, %r189, %r188; | |
| cvt.u64.u32 %rd628, %r190; | |
| $L__BB0_101: | |
| .loc 1 0 81 // bar.py:0:81 | |
| ld.param.b32 %r28, [fused_moe_kernel_param_18]; | |
| ld.param.b32 %r27, [fused_moe_kernel_param_17]; | |
| ld.param.b32 %r26, [fused_moe_kernel_param_16]; | |
| ld.param.b32 %r24, [fused_moe_kernel_param_14]; | |
| ld.param.b32 %r23, [fused_moe_kernel_param_13]; | |
| ld.param.b32 %r22, [fused_moe_kernel_param_12]; | |
| ld.param.b32 %r20, [fused_moe_kernel_param_9]; | |
| ld.param.b64 %rd117, [fused_moe_kernel_param_4]; | |
| ld.param.b64 %rd114, [fused_moe_kernel_param_1]; | |
| ld.param.b64 %rd113, [fused_moe_kernel_param_0]; | |
| and.b32 %r5, %r1, 127; | |
| cvt.s64.s32 %rd38, %r91; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd195, %rd71, %rd72; | |
| and.b64 %rd196, %rd195, -4294967296; | |
| setp.ne.b64 %p36, %rd196, 0; | |
| @%p36 bra $L__BB0_103; | |
| bra.uni $L__BB0_102; | |
| $L__BB0_103: | |
| rem.s64 %rd629, %rd71, %rd72; | |
| bra.uni $L__BB0_104; | |
| $L__BB0_102: | |
| cvt.u32.u64 %r191, %rd72; | |
| cvt.u32.u64 %r192, %rd71; | |
| rem.u32 %r193, %r192, %r191; | |
| cvt.u64.u32 %rd629, %r193; | |
| $L__BB0_104: | |
| .loc 1 0 81 // bar.py:0:81 | |
| ld.param.b64 %rd116, [fused_moe_kernel_param_3]; | |
| .loc 1 114 30 // bar.py:114:30 | |
| setp.lt.s32 %p37, %r89, %r21; | |
| .loc 1 135 31 // bar.py:135:31 | |
| shr.s32 %r203, %r89, 31; | |
| shr.u32 %r204, %r203, 29; | |
| add.s32 %r205, %r89, %r204; | |
| shr.s32 %r206, %r205, 3; | |
| .loc 1 135 39 // bar.py:135:39 | |
| mul.wide.s32 %rd73, %r206, %r22; | |
| .loc 1 135 58 // bar.py:135:58 | |
| and.b32 %r12, %r1, 1; | |
| shl.b32 %r207, %r12, 4; | |
| .loc 1 135 51 // bar.py:135:51 | |
| cvt.u64.u32 %rd74, %r207; | |
| .loc 1 135 8 // bar.py:135:8 | |
| add.s64 %rd201, %rd113, %rd73; | |
| add.s64 %rd197, %rd201, %rd74; | |
| .loc 1 139 24 // bar.py:139:24 | |
| cvt.s64.s32 %rd202, %r23; | |
| mul.lo.s64 %rd75, %rd38, %rd202; | |
| .loc 1 139 10 // bar.py:139:10 | |
| add.s64 %rd203, %rd114, %rd75; | |
| .loc 1 140 60 // bar.py:140:60 | |
| cvt.s64.s32 %rd204, %r24; | |
| mul.lo.s64 %rd76, %rd597, %rd204; | |
| .loc 1 140 11 // bar.py:140:11 | |
| add.s64 %rd205, %rd203, %rd76; | |
| add.s64 %rd198, %rd205, %rd74; | |
| .loc 1 151 56 // bar.py:151:56 | |
| shr.s32 %r208, %r90, 31; | |
| shr.u32 %r209, %r208, 29; | |
| add.s32 %r210, %r90, %r209; | |
| shr.s32 %r211, %r210, 3; | |
| .loc 1 151 65 // bar.py:151:65 | |
| mul.wide.s32 %rd206, %r211, %r26; | |
| .loc 1 151 41 // bar.py:151:41 | |
| shl.b64 %rd207, %rd206, 2; | |
| .loc 1 152 34 // bar.py:152:34 | |
| cvt.u32.u64 %r212, %rd598; | |
| shr.s32 %r213, %r212, 31; | |
| shr.u32 %r214, %r213, 25; | |
| add.s32 %r215, %r212, %r214; | |
| shr.s32 %r216, %r215, 7; | |
| cvt.u32.u64 %r217, %rd599; | |
| shr.s32 %r218, %r217, 31; | |
| shr.u32 %r219, %r218, 25; | |
| add.s32 %r220, %r217, %r219; | |
| shr.s32 %r221, %r220, 7; | |
| cvt.u32.u64 %r222, %rd600; | |
| shr.s32 %r223, %r222, 31; | |
| shr.u32 %r224, %r223, 25; | |
| add.s32 %r225, %r222, %r224; | |
| shr.s32 %r226, %r225, 7; | |
| cvt.u32.u64 %r227, %rd601; | |
| shr.s32 %r228, %r227, 31; | |
| shr.u32 %r229, %r228, 25; | |
| add.s32 %r230, %r227, %r229; | |
| shr.s32 %r231, %r230, 7; | |
| cvt.u32.u64 %r232, %rd602; | |
| shr.s32 %r233, %r232, 31; | |
| shr.u32 %r234, %r233, 25; | |
| add.s32 %r235, %r232, %r234; | |
| shr.s32 %r236, %r235, 7; | |
| cvt.u32.u64 %r237, %rd603; | |
| shr.s32 %r238, %r237, 31; | |
| shr.u32 %r239, %r238, 25; | |
| add.s32 %r240, %r237, %r239; | |
| shr.s32 %r241, %r240, 7; | |
| cvt.u32.u64 %r242, %rd604; | |
| shr.s32 %r243, %r242, 31; | |
| shr.u32 %r244, %r243, 25; | |
| add.s32 %r245, %r242, %r244; | |
| shr.s32 %r246, %r245, 7; | |
| cvt.u32.u64 %r247, %rd605; | |
| shr.s32 %r248, %r247, 31; | |
| shr.u32 %r249, %r248, 25; | |
| add.s32 %r250, %r247, %r249; | |
| shr.s32 %r251, %r250, 7; | |
| cvt.u32.u64 %r252, %rd606; | |
| shr.s32 %r253, %r252, 31; | |
| shr.u32 %r254, %r253, 25; | |
| add.s32 %r255, %r252, %r254; | |
| shr.s32 %r256, %r255, 7; | |
| cvt.u32.u64 %r257, %rd607; | |
| shr.s32 %r258, %r257, 31; | |
| shr.u32 %r259, %r258, 25; | |
| add.s32 %r260, %r257, %r259; | |
| shr.s32 %r261, %r260, 7; | |
| cvt.u32.u64 %r262, %rd608; | |
| shr.s32 %r263, %r262, 31; | |
| shr.u32 %r264, %r263, 25; | |
| add.s32 %r265, %r262, %r264; | |
| shr.s32 %r266, %r265, 7; | |
| cvt.u32.u64 %r267, %rd609; | |
| shr.s32 %r268, %r267, 31; | |
| shr.u32 %r269, %r268, 25; | |
| add.s32 %r270, %r267, %r269; | |
| shr.s32 %r271, %r270, 7; | |
| cvt.u32.u64 %r272, %rd610; | |
| shr.s32 %r273, %r272, 31; | |
| shr.u32 %r274, %r273, 25; | |
| add.s32 %r275, %r272, %r274; | |
| shr.s32 %r276, %r275, 7; | |
| cvt.u32.u64 %r277, %rd611; | |
| shr.s32 %r278, %r277, 31; | |
| shr.u32 %r279, %r278, 25; | |
| add.s32 %r280, %r277, %r279; | |
| shr.s32 %r281, %r280, 7; | |
| cvt.u32.u64 %r282, %rd612; | |
| shr.s32 %r283, %r282, 31; | |
| shr.u32 %r284, %r283, 25; | |
| add.s32 %r285, %r282, %r284; | |
| shr.s32 %r286, %r285, 7; | |
| cvt.u32.u64 %r287, %rd613; | |
| shr.s32 %r288, %r287, 31; | |
| shr.u32 %r289, %r288, 25; | |
| add.s32 %r290, %r287, %r289; | |
| shr.s32 %r291, %r290, 7; | |
| cvt.u32.u64 %r292, %rd614; | |
| shr.s32 %r293, %r292, 31; | |
| shr.u32 %r294, %r293, 25; | |
| add.s32 %r295, %r292, %r294; | |
| shr.s32 %r296, %r295, 7; | |
| cvt.u32.u64 %r297, %rd615; | |
| shr.s32 %r298, %r297, 31; | |
| shr.u32 %r299, %r298, 25; | |
| add.s32 %r300, %r297, %r299; | |
| shr.s32 %r301, %r300, 7; | |
| cvt.u32.u64 %r302, %rd616; | |
| shr.s32 %r303, %r302, 31; | |
| shr.u32 %r304, %r303, 25; | |
| add.s32 %r305, %r302, %r304; | |
| shr.s32 %r306, %r305, 7; | |
| cvt.u32.u64 %r307, %rd617; | |
| shr.s32 %r308, %r307, 31; | |
| shr.u32 %r309, %r308, 25; | |
| add.s32 %r310, %r307, %r309; | |
| shr.s32 %r311, %r310, 7; | |
| cvt.u32.u64 %r312, %rd618; | |
| shr.s32 %r313, %r312, 31; | |
| shr.u32 %r314, %r313, 25; | |
| add.s32 %r315, %r312, %r314; | |
| shr.s32 %r316, %r315, 7; | |
| cvt.u32.u64 %r317, %rd619; | |
| shr.s32 %r318, %r317, 31; | |
| shr.u32 %r319, %r318, 25; | |
| add.s32 %r320, %r317, %r319; | |
| shr.s32 %r321, %r320, 7; | |
| cvt.u32.u64 %r322, %rd620; | |
| shr.s32 %r323, %r322, 31; | |
| shr.u32 %r324, %r323, 25; | |
| add.s32 %r325, %r322, %r324; | |
| shr.s32 %r326, %r325, 7; | |
| cvt.u32.u64 %r327, %rd621; | |
| shr.s32 %r328, %r327, 31; | |
| shr.u32 %r329, %r328, 25; | |
| add.s32 %r330, %r327, %r329; | |
| shr.s32 %r331, %r330, 7; | |
| cvt.u32.u64 %r332, %rd622; | |
| shr.s32 %r333, %r332, 31; | |
| shr.u32 %r334, %r333, 25; | |
| add.s32 %r335, %r332, %r334; | |
| shr.s32 %r336, %r335, 7; | |
| cvt.u32.u64 %r337, %rd623; | |
| shr.s32 %r338, %r337, 31; | |
| shr.u32 %r339, %r338, 25; | |
| add.s32 %r340, %r337, %r339; | |
| shr.s32 %r341, %r340, 7; | |
| cvt.u32.u64 %r342, %rd624; | |
| shr.s32 %r343, %r342, 31; | |
| shr.u32 %r344, %r343, 25; | |
| add.s32 %r345, %r342, %r344; | |
| shr.s32 %r346, %r345, 7; | |
| cvt.u32.u64 %r347, %rd625; | |
| shr.s32 %r348, %r347, 31; | |
| shr.u32 %r349, %r348, 25; | |
| add.s32 %r350, %r347, %r349; | |
| shr.s32 %r351, %r350, 7; | |
| cvt.u32.u64 %r352, %rd626; | |
| shr.s32 %r353, %r352, 31; | |
| shr.u32 %r354, %r353, 25; | |
| add.s32 %r355, %r352, %r354; | |
| shr.s32 %r356, %r355, 7; | |
| cvt.u32.u64 %r357, %rd627; | |
| shr.s32 %r358, %r357, 31; | |
| shr.u32 %r359, %r358, 25; | |
| add.s32 %r360, %r357, %r359; | |
| shr.s32 %r361, %r360, 7; | |
| cvt.u32.u64 %r362, %rd628; | |
| shr.s32 %r363, %r362, 31; | |
| shr.u32 %r364, %r363, 25; | |
| add.s32 %r365, %r362, %r364; | |
| shr.s32 %r366, %r365, 7; | |
| cvt.u32.u64 %r367, %rd629; | |
| shr.s32 %r368, %r367, 31; | |
| shr.u32 %r369, %r368, 25; | |
| add.s32 %r370, %r367, %r369; | |
| shr.s32 %r371, %r370, 7; | |
| .loc 1 154 44 // bar.py:154:44 | |
| cvt.s64.s32 %rd208, %r27; | |
| mul.lo.s64 %rd209, %rd38, %rd208; | |
| .loc 1 154 30 // bar.py:154:30 | |
| shl.b64 %rd210, %rd209, 2; | |
| add.s64 %rd211, %rd117, %rd210; | |
| .loc 1 154 68 // bar.py:154:68 | |
| mul.wide.s32 %rd212, %r216, %r28; | |
| mul.wide.s32 %rd213, %r221, %r28; | |
| mul.wide.s32 %rd214, %r226, %r28; | |
| mul.wide.s32 %rd215, %r231, %r28; | |
| mul.wide.s32 %rd216, %r236, %r28; | |
| mul.wide.s32 %rd217, %r241, %r28; | |
| mul.wide.s32 %rd218, %r246, %r28; | |
| mul.wide.s32 %rd219, %r251, %r28; | |
| mul.wide.s32 %rd220, %r256, %r28; | |
| mul.wide.s32 %rd221, %r261, %r28; | |
| mul.wide.s32 %rd222, %r266, %r28; | |
| mul.wide.s32 %rd223, %r271, %r28; | |
| mul.wide.s32 %rd224, %r276, %r28; | |
| mul.wide.s32 %rd225, %r281, %r28; | |
| mul.wide.s32 %rd226, %r286, %r28; | |
| mul.wide.s32 %rd227, %r291, %r28; | |
| mul.wide.s32 %rd228, %r296, %r28; | |
| mul.wide.s32 %rd229, %r301, %r28; | |
| mul.wide.s32 %rd230, %r306, %r28; | |
| mul.wide.s32 %rd231, %r311, %r28; | |
| mul.wide.s32 %rd232, %r316, %r28; | |
| mul.wide.s32 %rd233, %r321, %r28; | |
| mul.wide.s32 %rd234, %r326, %r28; | |
| mul.wide.s32 %rd235, %r331, %r28; | |
| mul.wide.s32 %rd236, %r336, %r28; | |
| mul.wide.s32 %rd237, %r341, %r28; | |
| mul.wide.s32 %rd238, %r346, %r28; | |
| mul.wide.s32 %rd239, %r351, %r28; | |
| mul.wide.s32 %rd240, %r356, %r28; | |
| mul.wide.s32 %rd241, %r361, %r28; | |
| mul.wide.s32 %rd242, %r366, %r28; | |
| mul.wide.s32 %rd243, %r371, %r28; | |
| .loc 1 154 57 // bar.py:154:57 | |
| shl.b64 %rd244, %rd212, 2; | |
| shl.b64 %rd245, %rd213, 2; | |
| shl.b64 %rd246, %rd214, 2; | |
| shl.b64 %rd247, %rd215, 2; | |
| shl.b64 %rd248, %rd216, 2; | |
| shl.b64 %rd249, %rd217, 2; | |
| shl.b64 %rd250, %rd218, 2; | |
| shl.b64 %rd251, %rd219, 2; | |
| shl.b64 %rd252, %rd220, 2; | |
| shl.b64 %rd253, %rd221, 2; | |
| shl.b64 %rd254, %rd222, 2; | |
| shl.b64 %rd255, %rd223, 2; | |
| shl.b64 %rd256, %rd224, 2; | |
| shl.b64 %rd257, %rd225, 2; | |
| shl.b64 %rd258, %rd226, 2; | |
| shl.b64 %rd259, %rd227, 2; | |
| shl.b64 %rd260, %rd228, 2; | |
| shl.b64 %rd261, %rd229, 2; | |
| shl.b64 %rd262, %rd230, 2; | |
| shl.b64 %rd263, %rd231, 2; | |
| shl.b64 %rd264, %rd232, 2; | |
| shl.b64 %rd265, %rd233, 2; | |
| shl.b64 %rd266, %rd234, 2; | |
| shl.b64 %rd267, %rd235, 2; | |
| shl.b64 %rd268, %rd236, 2; | |
| shl.b64 %rd269, %rd237, 2; | |
| shl.b64 %rd270, %rd238, 2; | |
| shl.b64 %rd271, %rd239, 2; | |
| shl.b64 %rd272, %rd240, 2; | |
| shl.b64 %rd273, %rd241, 2; | |
| shl.b64 %rd274, %rd242, 2; | |
| shl.b64 %rd275, %rd243, 2; | |
| $L__tmp6: | |
| .loc 2 43 17 // standard.py:43:17 @[ bar.py:172:33 ] | |
| add.s32 %r13, %r20, 31; | |
| $L__tmp7: | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.eq.b32 %p60, %r5, 0; | |
| add.s32 %r671, %r29, 12544; | |
| // begin inline asm | |
| @%p60 mbarrier.init.shared::cta.b64 [%r671], 1; | |
| // end inline asm | |
| bar.sync 0; | |
| add.s32 %r194, %r29, 12552; | |
| // begin inline asm | |
| @%p60 mbarrier.init.shared::cta.b64 [%r194], 1; | |
| // end inline asm | |
| setp.lt.s32 %p38, %r13, 32; | |
| setp.gt.s32 %p39, %r13, 31; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p40, %r207, %r20; | |
| .loc 1 174 12 // bar.py:174:12 | |
| shl.b32 %r15, %r1, 4; | |
| and.b32 %r375, %r15, 1904; | |
| bfe.s32 %r376, %r1, 3, 1; | |
| and.b32 %r377, %r376, 144; | |
| xor.b32 %r378, %r377, %r375; | |
| add.s32 %r195, %r29, %r378; | |
| selp.b32 %r379, 16, 0, %p37; | |
| selp.b32 %r380, %r379, 0, %p40; | |
| selp.b32 %r196, %r380, 0, %p39; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r195 + 0 ], [ %rd197 + 0 ], 0x10, %r196; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r197, %r195, 6144; | |
| selp.b32 %r381, 16, 0, %p39; | |
| selp.b32 %r198, %r381, 0, %p40; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r197 + 0 ], [ %rd198 + 0 ], 0x10, %r198; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.gt.s32 %p41, %r13, 63; | |
| .loc 1 197 18 // bar.py:197:18 | |
| add.s64 %rd199, %rd197, 32; | |
| .loc 1 198 18 // bar.py:198:18 | |
| add.s64 %rd200, %rd198, 32; | |
| .loc 1 175 62 // bar.py:175:62 | |
| add.s32 %r382, %r20, -32; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p42, %r207, %r382; | |
| .loc 1 174 12 // bar.py:174:12 | |
| bar.sync 0; | |
| add.s32 %r199, %r195, 2048; | |
| selp.b32 %r383, %r379, 0, %p42; | |
| selp.b32 %r200, %r383, 0, %p41; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r199 + 0 ], [ %rd199 + 0 ], 0x10, %r200; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r201, %r195, 8192; | |
| selp.b32 %r384, 16, 0, %p42; | |
| selp.b32 %r202, %r384, 0, %p41; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r201 + 0 ], [ %rd200 + 0 ], 0x10, %r202; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 174 12 // bar.py:174:12 | |
| cp.async.wait_group 2; | |
| bar.sync 0; | |
| .loc 1 189 41 // bar.py:189:41 | |
| shr.u32 %r385, %r1, 5; | |
| shfl.sync.idx.b32 %r16, %r385, 0, 31, -1; | |
| setp.ne.b32 %p43, %r16, 0; | |
| or.pred %p44, %p38, %p43; | |
| @%p44 bra $L__BB0_106; | |
| // %bb.105: | |
| elect.sync %r387|%p46, -1; | |
| bfe.u32 %r388, %r29, 4, 14; | |
| cvt.u64.u32 %rd279, %r388; | |
| add.s32 %r389, %r29, 6144; | |
| bfe.u32 %r390, %r389, 4, 14; | |
| cvt.u64.u32 %rd280, %r390; | |
| or.b64 %rd276, %rd279, -4611615580963733504; | |
| or.b64 %rd277, %rd280, -4611615580963733504; | |
| mov.b32 %r386, 68157456; | |
| mov.pred %p45, 0; | |
| // begin inline asm | |
| @%p46 tcgen05.mma.cta_group::1.kind::f8f6f4 [ %r2 + 0 ], %rd276, %rd277, %r386, %p45; | |
| // end inline asm | |
| add.s32 %r391, %r29, 12544; | |
| cvt.u64.u32 %rd278, %r391; | |
| // begin inline asm | |
| @%p46 tcgen05.commit.cta_group::1.mbarrier::arrive::one.b64 [%rd278]; | |
| // end inline asm | |
| $L__BB0_106: | |
| .loc 1 0 41 // bar.py:0:41 | |
| add.s64 %rd77, %rd116, %rd207; | |
| add.s64 %rd78, %rd211, %rd244; | |
| add.s64 %rd79, %rd211, %rd245; | |
| add.s64 %rd80, %rd211, %rd246; | |
| add.s64 %rd81, %rd211, %rd247; | |
| add.s64 %rd82, %rd211, %rd248; | |
| add.s64 %rd83, %rd211, %rd249; | |
| add.s64 %rd84, %rd211, %rd250; | |
| add.s64 %rd85, %rd211, %rd251; | |
| add.s64 %rd86, %rd211, %rd252; | |
| add.s64 %rd87, %rd211, %rd253; | |
| add.s64 %rd88, %rd211, %rd254; | |
| add.s64 %rd89, %rd211, %rd255; | |
| add.s64 %rd90, %rd211, %rd256; | |
| add.s64 %rd91, %rd211, %rd257; | |
| add.s64 %rd92, %rd211, %rd258; | |
| add.s64 %rd93, %rd211, %rd259; | |
| add.s64 %rd94, %rd211, %rd260; | |
| add.s64 %rd95, %rd211, %rd261; | |
| add.s64 %rd96, %rd211, %rd262; | |
| add.s64 %rd97, %rd211, %rd263; | |
| add.s64 %rd98, %rd211, %rd264; | |
| add.s64 %rd99, %rd211, %rd265; | |
| add.s64 %rd100, %rd211, %rd266; | |
| add.s64 %rd101, %rd211, %rd267; | |
| add.s64 %rd102, %rd211, %rd268; | |
| add.s64 %rd103, %rd211, %rd269; | |
| add.s64 %rd104, %rd211, %rd270; | |
| add.s64 %rd105, %rd211, %rd271; | |
| add.s64 %rd106, %rd211, %rd272; | |
| add.s64 %rd107, %rd211, %rd273; | |
| add.s64 %rd108, %rd211, %rd274; | |
| add.s64 %rd109, %rd211, %rd275; | |
| cvt.u32.u64 %r396, %rd74; | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.gt.s32 %p47, %r13, 95; | |
| setp.lt.s32 %p48, %r13, 64; | |
| .loc 1 197 18 // bar.py:197:18 | |
| add.s64 %rd281, %rd197, 64; | |
| .loc 1 198 18 // bar.py:198:18 | |
| add.s64 %rd282, %rd198, 64; | |
| .loc 1 175 62 // bar.py:175:62 | |
| add.s32 %r397, %r20, -64; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p49, %r396, %r397; | |
| .loc 1 174 12 // bar.py:174:12 | |
| bar.sync 0; | |
| add.s32 %r392, %r195, 4096; | |
| selp.b32 %r398, %r379, 0, %p49; | |
| selp.b32 %r393, %r398, 0, %p47; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r392 + 0 ], [ %rd281 + 0 ], 0x10, %r393; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r394, %r195, 10240; | |
| selp.b32 %r399, 16, 0, %p49; | |
| selp.b32 %r395, %r399, 0, %p47; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r394 + 0 ], [ %rd282 + 0 ], 0x10, %r395; | |
| // end inline asm | |
| cp.async.commit_group; | |
| mov.b32 %r400, 0f00000000; | |
| mov.b64 %rd632, {%r400, %r400}; | |
| mov.b32 %r672, 0; | |
| mov.b64 %rd631, 0; | |
| setp.eq.b32 %p79, %r11, 0; | |
| setp.lt.u32 %p80, %r5, 32; | |
| shl.b32 %r661, %r9, 3; | |
| shl.b32 %r662, %r16, 21; | |
| setp.lt.s32 %p81, %r90, %r21; | |
| mov.b64 %rd633, %rd632; | |
| mov.b64 %rd634, %rd632; | |
| mov.b64 %rd635, %rd632; | |
| mov.b64 %rd636, %rd632; | |
| mov.b64 %rd637, %rd632; | |
| mov.b64 %rd638, %rd632; | |
| mov.b64 %rd639, %rd632; | |
| mov.b64 %rd640, %rd632; | |
| mov.b64 %rd641, %rd632; | |
| mov.b64 %rd642, %rd632; | |
| mov.b64 %rd643, %rd632; | |
| mov.b64 %rd644, %rd632; | |
| mov.b64 %rd645, %rd632; | |
| mov.b64 %rd646, %rd632; | |
| mov.b64 %rd647, %rd632; | |
| .loc 1 172 22 // bar.py:172:22 | |
| @%p48 bra $L__BB0_112; | |
| // %bb.107: // %.lr.ph | |
| .loc 1 0 22 // bar.py:0:22 | |
| shr.s32 %r372, %r13, 31; | |
| shr.u32 %r373, %r372, 27; | |
| add.s32 %r374, %r13, %r373; | |
| shr.s32 %r14, %r374, 5; | |
| add.s32 %r17, %r14, -3; | |
| and.pred %p50, %p79, %p80; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s32 %r402, %r29, %r661; | |
| add.s32 %r406, %r402, 12288; | |
| and.b32 %r403, %r662, 6291456; | |
| add.s32 %r470, %r403, %r2; | |
| max.s32 %r404, %r14, 2; | |
| add.s32 %r407, %r402, 12292; | |
| add.s32 %r408, %r402, 12296; | |
| add.s32 %r409, %r402, 12300; | |
| add.s32 %r410, %r402, 12304; | |
| add.s32 %r411, %r402, 12308; | |
| add.s32 %r412, %r402, 12312; | |
| add.s32 %r413, %r402, 12316; | |
| add.s32 %r414, %r402, 12320; | |
| add.s32 %r415, %r402, 12324; | |
| add.s32 %r416, %r402, 12328; | |
| add.s32 %r417, %r402, 12332; | |
| add.s32 %r418, %r402, 12336; | |
| add.s32 %r419, %r402, 12340; | |
| add.s32 %r420, %r402, 12344; | |
| add.s32 %r421, %r402, 12348; | |
| add.s32 %r422, %r402, 12352; | |
| add.s32 %r423, %r402, 12356; | |
| add.s32 %r424, %r402, 12360; | |
| add.s32 %r425, %r402, 12364; | |
| add.s32 %r426, %r402, 12368; | |
| add.s32 %r427, %r402, 12372; | |
| add.s32 %r428, %r402, 12376; | |
| add.s32 %r429, %r402, 12380; | |
| add.s32 %r430, %r402, 12384; | |
| add.s32 %r431, %r402, 12388; | |
| add.s32 %r432, %r402, 12392; | |
| add.s32 %r433, %r402, 12396; | |
| add.s32 %r434, %r402, 12400; | |
| add.s32 %r435, %r402, 12404; | |
| add.s32 %r436, %r402, 12408; | |
| add.s32 %r437, %r402, 12412; | |
| add.s64 %rd283, %rd76, %rd75; | |
| add.s64 %rd284, %rd283, %rd74; | |
| add.s64 %rd285, %rd284, %rd114; | |
| add.s64 %rd110, %rd285, 96; | |
| add.s64 %rd286, %rd73, %rd74; | |
| add.s64 %rd287, %rd286, %rd113; | |
| add.s64 %rd111, %rd287, 96; | |
| mad.wide.u32 %rd112, %r404, 32, -32; | |
| mov.b64 %rd632, {%r400, %r400}; | |
| add.s32 %r671, %r29, 12544; | |
| add.s32 %r664, %r20, -96; | |
| mov.b32 %r668, 1; | |
| mov.b32 %r667, 2; | |
| mov.b32 %r401, 0; | |
| mov.b64 %rd630, 0; | |
| mov.b32 %r665, %r401; | |
| mov.b32 %r666, %r401; | |
| mov.b32 %r669, %r401; | |
| mov.b32 %r670, %r401; | |
| mov.b64 %rd633, %rd632; | |
| mov.b64 %rd634, %rd632; | |
| mov.b64 %rd635, %rd632; | |
| mov.b64 %rd636, %rd632; | |
| mov.b64 %rd637, %rd632; | |
| mov.b64 %rd638, %rd632; | |
| mov.b64 %rd639, %rd632; | |
| mov.b64 %rd640, %rd632; | |
| mov.b64 %rd641, %rd632; | |
| mov.b64 %rd642, %rd632; | |
| mov.b64 %rd643, %rd632; | |
| mov.b64 %rd644, %rd632; | |
| mov.b64 %rd645, %rd632; | |
| mov.b64 %rd646, %rd632; | |
| mov.b64 %rd647, %rd632; | |
| bra.uni $L__BB0_108; | |
| $L__BB0_110: // in Loop: Header=BB0_108 Depth=1 | |
| .loc 1 0 0 // bar.py:0 | |
| fma.rn.f32x2 %rd647, %rd337, %rd420, %rd647; | |
| fma.rn.f32x2 %rd646, %rd336, %rd419, %rd646; | |
| fma.rn.f32x2 %rd645, %rd335, %rd418, %rd645; | |
| fma.rn.f32x2 %rd644, %rd334, %rd417, %rd644; | |
| fma.rn.f32x2 %rd643, %rd333, %rd416, %rd643; | |
| fma.rn.f32x2 %rd642, %rd332, %rd415, %rd642; | |
| fma.rn.f32x2 %rd641, %rd331, %rd414, %rd641; | |
| fma.rn.f32x2 %rd640, %rd330, %rd413, %rd640; | |
| fma.rn.f32x2 %rd639, %rd329, %rd412, %rd639; | |
| fma.rn.f32x2 %rd638, %rd328, %rd411, %rd638; | |
| fma.rn.f32x2 %rd637, %rd327, %rd410, %rd637; | |
| fma.rn.f32x2 %rd636, %rd326, %rd409, %rd636; | |
| fma.rn.f32x2 %rd635, %rd325, %rd408, %rd635; | |
| fma.rn.f32x2 %rd634, %rd324, %rd407, %rd634; | |
| fma.rn.f32x2 %rd633, %rd323, %rd406, %rd633; | |
| fma.rn.f32x2 %rd632, %rd322, %rd405, %rd632; | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.lt.s32 %p54, %r670, %r17; | |
| .loc 1 189 41 // bar.py:189:41 | |
| add.s32 %r487, %r668, 1; | |
| setp.gt.s32 %p55, %r487, 1; | |
| selp.b32 %r668, 0, %r487, %p55; | |
| selp.b32 %r488, 1, 0, %p55; | |
| xor.b32 %r669, %r672, %r488; | |
| .loc 1 198 18 // bar.py:198:18 | |
| add.s64 %rd426, %rd111, %rd630; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s64 %rd427, %rd110, %rd630; | |
| add.s32 %r489, %r667, 1; | |
| setp.gt.s32 %p56, %r489, 2; | |
| selp.b32 %r667, 0, %r489, %p56; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p57, %r396, %r664; | |
| .loc 1 174 12 // bar.py:174:12 | |
| shl.b32 %r490, %r667, 11; | |
| bar.sync 0; | |
| add.s32 %r483, %r195, %r490; | |
| selp.b32 %r491, 16, 0, %p57; | |
| selp.b32 %r492, %r491, 0, %p37; | |
| selp.b32 %r484, %r492, 0, %p54; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r483 + 0 ], [ %rd426 + 0 ], 0x10, %r484; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r485, %r483, 6144; | |
| selp.b32 %r486, %r491, 0, %p54; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r485 + 0 ], [ %rd427 + 0 ], 0x10, %r486; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s32 %r670, %r670, 1; | |
| add.s64 %rd630, %rd630, 32; | |
| cvt.u32.u64 %r18, %rd630; | |
| cvt.u32.u64 %r493, %rd112; | |
| add.s32 %r664, %r664, -32; | |
| setp.ne.b32 %p58, %r493, %r18; | |
| mov.b32 %r665, %r672; | |
| @%p58 bra $L__BB0_108; | |
| bra.uni $L__BB0_111; | |
| $L__BB0_108: // =>This Inner Loop Header: Depth=1 | |
| .loc 1 0 22 // bar.py:0:22 | |
| mov.b32 %r672, %r669; | |
| cvt.u32.u64 %r471, %rd630; | |
| .loc 1 184 37 // bar.py:184:37 | |
| shr.u32 %r472, %r471, 7; | |
| .loc 1 186 35 // bar.py:186:35 | |
| mul.wide.u32 %rd321, %r472, 4; | |
| add.s64 %rd288, %rd77, %rd321; | |
| .loc 1 186 20 // bar.py:186:20 | |
| // begin inline asm | |
| mov.u32 %r405, %r401; | |
| @%p81 ld.global.b32 { %r405 }, [ %rd288 + 0 ]; | |
| // end inline asm | |
| .loc 1 188 49 // bar.py:188:49 | |
| add.s64 %rd289, %rd78, %rd321; | |
| add.s64 %rd290, %rd79, %rd321; | |
| add.s64 %rd291, %rd80, %rd321; | |
| add.s64 %rd292, %rd81, %rd321; | |
| add.s64 %rd293, %rd82, %rd321; | |
| add.s64 %rd294, %rd83, %rd321; | |
| add.s64 %rd295, %rd84, %rd321; | |
| add.s64 %rd296, %rd85, %rd321; | |
| add.s64 %rd297, %rd86, %rd321; | |
| add.s64 %rd298, %rd87, %rd321; | |
| add.s64 %rd299, %rd88, %rd321; | |
| add.s64 %rd300, %rd89, %rd321; | |
| add.s64 %rd301, %rd90, %rd321; | |
| add.s64 %rd302, %rd91, %rd321; | |
| add.s64 %rd303, %rd92, %rd321; | |
| add.s64 %rd304, %rd93, %rd321; | |
| add.s64 %rd305, %rd94, %rd321; | |
| add.s64 %rd306, %rd95, %rd321; | |
| add.s64 %rd307, %rd96, %rd321; | |
| add.s64 %rd308, %rd97, %rd321; | |
| add.s64 %rd309, %rd98, %rd321; | |
| add.s64 %rd310, %rd99, %rd321; | |
| add.s64 %rd311, %rd100, %rd321; | |
| add.s64 %rd312, %rd101, %rd321; | |
| add.s64 %rd313, %rd102, %rd321; | |
| add.s64 %rd314, %rd103, %rd321; | |
| add.s64 %rd315, %rd104, %rd321; | |
| add.s64 %rd316, %rd105, %rd321; | |
| add.s64 %rd317, %rd106, %rd321; | |
| add.s64 %rd318, %rd107, %rd321; | |
| add.s64 %rd319, %rd108, %rd321; | |
| add.s64 %rd320, %rd109, %rd321; | |
| .loc 1 188 34 // bar.py:188:34 | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r406 + 0 ], [ %rd289 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r407 + 0 ], [ %rd290 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r408 + 0 ], [ %rd291 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r409 + 0 ], [ %rd292 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r410 + 0 ], [ %rd293 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r411 + 0 ], [ %rd294 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r412 + 0 ], [ %rd295 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r413 + 0 ], [ %rd296 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r414 + 0 ], [ %rd297 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r415 + 0 ], [ %rd298 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r416 + 0 ], [ %rd299 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r417 + 0 ], [ %rd300 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r418 + 0 ], [ %rd301 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r419 + 0 ], [ %rd302 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r420 + 0 ], [ %rd303 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r421 + 0 ], [ %rd304 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r422 + 0 ], [ %rd305 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r423 + 0 ], [ %rd306 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r424 + 0 ], [ %rd307 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r425 + 0 ], [ %rd308 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r426 + 0 ], [ %rd309 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r427 + 0 ], [ %rd310 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r428 + 0 ], [ %rd311 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r429 + 0 ], [ %rd312 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r430 + 0 ], [ %rd313 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r431 + 0 ], [ %rd314 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r432 + 0 ], [ %rd315 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r433 + 0 ], [ %rd316 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r434 + 0 ], [ %rd317 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r435 + 0 ], [ %rd318 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r436 + 0 ], [ %rd319 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 cp.async.ca.shared.global [ %r437 + 0 ], [ %rd320 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| { | |
| .reg .pred complete; | |
| waitLoop: | |
| mbarrier.try_wait.parity.shared.b64 complete, [%r671], %r665; | |
| @!complete bra.uni waitLoop; | |
| } | |
| // end inline asm | |
| .loc 1 188 34 // bar.py:188:34 | |
| cp.async.wait_group 0; | |
| bar.sync 0; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s32 %r473, %r666, 1; | |
| setp.gt.s32 %p51, %r473, 2; | |
| selp.b32 %r666, 0, %r473, %p51; | |
| shl.b32 %r474, %r668, 3; | |
| add.s32 %r475, %r29, %r474; | |
| add.s32 %r671, %r475, 12544; | |
| .loc 1 188 34 // bar.py:188:34 | |
| ld.shared.v2.b64 {%rd322, %rd323}, [%r406]; | |
| ld.shared.v2.b64 {%rd324, %rd325}, [%r406+16]; | |
| ld.shared.v2.b64 {%rd326, %rd327}, [%r406+32]; | |
| ld.shared.v2.b64 {%rd328, %rd329}, [%r406+48]; | |
| ld.shared.v2.b64 {%rd330, %rd331}, [%r406+64]; | |
| ld.shared.v2.b64 {%rd332, %rd333}, [%r406+80]; | |
| ld.shared.v2.b64 {%rd334, %rd335}, [%r406+96]; | |
| ld.shared.v2.b64 {%rd336, %rd337}, [%r406+112]; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| tcgen05.ld.sync.aligned.16x32bx2.x32.b32 {%r438, %r439, %r440, %r441, %r442, %r443, %r444, %r445, %r446, %r447, %r448, %r449, %r450, %r451, %r452, %r453, %r454, %r455, %r456, %r457, %r458, %r459, %r460, %r461, %r462, %r463, %r464, %r465, %r466, %r467, %r468, %r469}, [%r470 + 0], 32; | |
| // end inline asm | |
| cvt.u64.u32 %rd338, %r468; | |
| cvt.u64.u32 %rd339, %r469; | |
| shl.b64 %rd340, %rd339, 32; | |
| or.b64 %rd341, %rd338, %rd340; | |
| cvt.u64.u32 %rd342, %r466; | |
| cvt.u64.u32 %rd343, %r467; | |
| shl.b64 %rd344, %rd343, 32; | |
| or.b64 %rd345, %rd342, %rd344; | |
| cvt.u64.u32 %rd346, %r464; | |
| cvt.u64.u32 %rd347, %r465; | |
| shl.b64 %rd348, %rd347, 32; | |
| or.b64 %rd349, %rd346, %rd348; | |
| cvt.u64.u32 %rd350, %r462; | |
| cvt.u64.u32 %rd351, %r463; | |
| shl.b64 %rd352, %rd351, 32; | |
| or.b64 %rd353, %rd350, %rd352; | |
| cvt.u64.u32 %rd354, %r460; | |
| cvt.u64.u32 %rd355, %r461; | |
| shl.b64 %rd356, %rd355, 32; | |
| or.b64 %rd357, %rd354, %rd356; | |
| cvt.u64.u32 %rd358, %r458; | |
| cvt.u64.u32 %rd359, %r459; | |
| shl.b64 %rd360, %rd359, 32; | |
| or.b64 %rd361, %rd358, %rd360; | |
| cvt.u64.u32 %rd362, %r456; | |
| cvt.u64.u32 %rd363, %r457; | |
| shl.b64 %rd364, %rd363, 32; | |
| or.b64 %rd365, %rd362, %rd364; | |
| cvt.u64.u32 %rd366, %r454; | |
| cvt.u64.u32 %rd367, %r455; | |
| shl.b64 %rd368, %rd367, 32; | |
| or.b64 %rd369, %rd366, %rd368; | |
| cvt.u64.u32 %rd370, %r452; | |
| cvt.u64.u32 %rd371, %r453; | |
| shl.b64 %rd372, %rd371, 32; | |
| or.b64 %rd373, %rd370, %rd372; | |
| cvt.u64.u32 %rd374, %r450; | |
| cvt.u64.u32 %rd375, %r451; | |
| shl.b64 %rd376, %rd375, 32; | |
| or.b64 %rd377, %rd374, %rd376; | |
| cvt.u64.u32 %rd378, %r448; | |
| cvt.u64.u32 %rd379, %r449; | |
| shl.b64 %rd380, %rd379, 32; | |
| or.b64 %rd381, %rd378, %rd380; | |
| cvt.u64.u32 %rd382, %r446; | |
| cvt.u64.u32 %rd383, %r447; | |
| shl.b64 %rd384, %rd383, 32; | |
| or.b64 %rd385, %rd382, %rd384; | |
| cvt.u64.u32 %rd386, %r444; | |
| cvt.u64.u32 %rd387, %r445; | |
| shl.b64 %rd388, %rd387, 32; | |
| or.b64 %rd389, %rd386, %rd388; | |
| cvt.u64.u32 %rd390, %r442; | |
| cvt.u64.u32 %rd391, %r443; | |
| shl.b64 %rd392, %rd391, 32; | |
| or.b64 %rd393, %rd390, %rd392; | |
| cvt.u64.u32 %rd394, %r440; | |
| cvt.u64.u32 %rd395, %r441; | |
| shl.b64 %rd396, %rd395, 32; | |
| or.b64 %rd397, %rd394, %rd396; | |
| cvt.u64.u32 %rd398, %r438; | |
| cvt.u64.u32 %rd399, %r439; | |
| shl.b64 %rd400, %rd399, 32; | |
| or.b64 %rd401, %rd398, %rd400; | |
| tcgen05.wait::ld.sync.aligned; | |
| .loc 1 189 46 // bar.py:189:46 | |
| cvt.u64.u32 %rd402, %r405; | |
| shl.b64 %rd403, %rd402, 32; | |
| or.b64 %rd404, %rd402, %rd403; | |
| mul.f32x2 %rd405, %rd404, %rd401; | |
| mul.f32x2 %rd406, %rd404, %rd397; | |
| mul.f32x2 %rd407, %rd404, %rd393; | |
| mul.f32x2 %rd408, %rd404, %rd389; | |
| mul.f32x2 %rd409, %rd404, %rd385; | |
| mul.f32x2 %rd410, %rd404, %rd381; | |
| mul.f32x2 %rd411, %rd404, %rd377; | |
| mul.f32x2 %rd412, %rd404, %rd373; | |
| mul.f32x2 %rd413, %rd404, %rd369; | |
| mul.f32x2 %rd414, %rd404, %rd365; | |
| mul.f32x2 %rd415, %rd404, %rd361; | |
| mul.f32x2 %rd416, %rd404, %rd357; | |
| mul.f32x2 %rd417, %rd404, %rd353; | |
| mul.f32x2 %rd418, %rd404, %rd349; | |
| mul.f32x2 %rd419, %rd404, %rd345; | |
| mul.f32x2 %rd420, %rd404, %rd341; | |
| .loc 1 189 41 // bar.py:189:41 | |
| @%p43 bra $L__BB0_110; | |
| // %bb.109: // in Loop: Header=BB0_108 Depth=1 | |
| .loc 1 174 12 // bar.py:174:12 | |
| shl.b32 %r477, %r666, 11; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r478, %r29, %r477; | |
| add.s32 %r479, %r478, 6144; | |
| .loc 1 189 41 // bar.py:189:41 | |
| elect.sync %r480|%p53, -1; | |
| bfe.u32 %r481, %r478, 4, 14; | |
| cvt.u64.u32 %rd424, %r481; | |
| bfe.u32 %r482, %r479, 4, 14; | |
| cvt.u64.u32 %rd425, %r482; | |
| or.b64 %rd421, %rd424, -4611615580963733504; | |
| or.b64 %rd422, %rd425, -4611615580963733504; | |
| mov.b32 %r476, 68157456; | |
| mov.pred %p52, 0; | |
| // begin inline asm | |
| @%p53 tcgen05.mma.cta_group::1.kind::f8f6f4 [ %r2 + 0 ], %rd421, %rd422, %r476, %p52; | |
| // end inline asm | |
| cvt.u64.u32 %rd423, %r671; | |
| // begin inline asm | |
| @%p53 tcgen05.commit.cta_group::1.mbarrier::arrive::one.b64 [%rd423]; | |
| // end inline asm | |
| bra.uni $L__BB0_110; | |
| $L__BB0_111: // %._crit_edge.loopexit | |
| .loc 1 184 37 // bar.py:184:37 | |
| shr.s32 %r494, %r18, 31; | |
| shr.u32 %r495, %r494, 25; | |
| add.s32 %r496, %r18, %r495; | |
| shr.s32 %r497, %r496, 7; | |
| .loc 1 186 35 // bar.py:186:35 | |
| cvt.s64.s32 %rd631, %r497; | |
| $L__BB0_112: // %._crit_edge | |
| .loc 1 172 22 // bar.py:172:22 | |
| @%p38 bra $L__BB0_114; | |
| // %bb.113: | |
| .loc 1 186 35 // bar.py:186:35 | |
| shl.b64 %rd461, %rd631, 2; | |
| add.s64 %rd428, %rd77, %rd461; | |
| mov.b32 %r499, 0; | |
| .loc 1 186 20 // bar.py:186:20 | |
| // begin inline asm | |
| mov.u32 %r498, %r499; | |
| @%p81 ld.global.b32 { %r498 }, [ %rd428 + 0 ]; | |
| // end inline asm | |
| .loc 1 188 49 // bar.py:188:49 | |
| add.s64 %rd429, %rd78, %rd461; | |
| add.s64 %rd430, %rd79, %rd461; | |
| add.s64 %rd431, %rd80, %rd461; | |
| add.s64 %rd432, %rd81, %rd461; | |
| add.s64 %rd433, %rd82, %rd461; | |
| add.s64 %rd434, %rd83, %rd461; | |
| add.s64 %rd435, %rd84, %rd461; | |
| add.s64 %rd436, %rd85, %rd461; | |
| add.s64 %rd437, %rd86, %rd461; | |
| add.s64 %rd438, %rd87, %rd461; | |
| add.s64 %rd439, %rd88, %rd461; | |
| add.s64 %rd440, %rd89, %rd461; | |
| add.s64 %rd441, %rd90, %rd461; | |
| add.s64 %rd442, %rd91, %rd461; | |
| add.s64 %rd443, %rd92, %rd461; | |
| add.s64 %rd444, %rd93, %rd461; | |
| add.s64 %rd445, %rd94, %rd461; | |
| add.s64 %rd446, %rd95, %rd461; | |
| add.s64 %rd447, %rd96, %rd461; | |
| add.s64 %rd448, %rd97, %rd461; | |
| add.s64 %rd449, %rd98, %rd461; | |
| add.s64 %rd450, %rd99, %rd461; | |
| add.s64 %rd451, %rd100, %rd461; | |
| add.s64 %rd452, %rd101, %rd461; | |
| add.s64 %rd453, %rd102, %rd461; | |
| add.s64 %rd454, %rd103, %rd461; | |
| add.s64 %rd455, %rd104, %rd461; | |
| add.s64 %rd456, %rd105, %rd461; | |
| add.s64 %rd457, %rd106, %rd461; | |
| add.s64 %rd458, %rd107, %rd461; | |
| add.s64 %rd459, %rd108, %rd461; | |
| add.s64 %rd460, %rd109, %rd461; | |
| .loc 1 188 34 // bar.py:188:34 | |
| and.pred %p59, %p79, %p80; | |
| add.s32 %r565, %r29, %r661; | |
| add.s32 %r500, %r565, 12288; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r500 + 0 ], [ %rd429 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r501, %r565, 12292; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r501 + 0 ], [ %rd430 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r502, %r565, 12296; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r502 + 0 ], [ %rd431 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r503, %r565, 12300; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r503 + 0 ], [ %rd432 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r504, %r565, 12304; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r504 + 0 ], [ %rd433 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r505, %r565, 12308; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r505 + 0 ], [ %rd434 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r506, %r565, 12312; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r506 + 0 ], [ %rd435 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r507, %r565, 12316; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r507 + 0 ], [ %rd436 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r508, %r565, 12320; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r508 + 0 ], [ %rd437 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r509, %r565, 12324; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r509 + 0 ], [ %rd438 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r510, %r565, 12328; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r510 + 0 ], [ %rd439 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r511, %r565, 12332; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r511 + 0 ], [ %rd440 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r512, %r565, 12336; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r512 + 0 ], [ %rd441 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r513, %r565, 12340; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r513 + 0 ], [ %rd442 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r514, %r565, 12344; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r514 + 0 ], [ %rd443 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r515, %r565, 12348; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r515 + 0 ], [ %rd444 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r516, %r565, 12352; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r516 + 0 ], [ %rd445 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r517, %r565, 12356; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r517 + 0 ], [ %rd446 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r518, %r565, 12360; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r518 + 0 ], [ %rd447 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r519, %r565, 12364; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r519 + 0 ], [ %rd448 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r520, %r565, 12368; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r520 + 0 ], [ %rd449 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r521, %r565, 12372; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r521 + 0 ], [ %rd450 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r522, %r565, 12376; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r522 + 0 ], [ %rd451 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r523, %r565, 12380; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r523 + 0 ], [ %rd452 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r524, %r565, 12384; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r524 + 0 ], [ %rd453 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r525, %r565, 12388; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r525 + 0 ], [ %rd454 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r526, %r565, 12392; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r526 + 0 ], [ %rd455 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r527, %r565, 12396; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r527 + 0 ], [ %rd456 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r528, %r565, 12400; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r528 + 0 ], [ %rd457 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r529, %r565, 12404; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r529 + 0 ], [ %rd458 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r530, %r565, 12408; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r530 + 0 ], [ %rd459 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| add.s32 %r531, %r565, 12412; | |
| // begin inline asm | |
| @%p59 cp.async.ca.shared.global [ %r531 + 0 ], [ %rd460 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| { | |
| .reg .pred complete; | |
| waitLoop: | |
| mbarrier.try_wait.parity.shared.b64 complete, [%r671], %r672; | |
| @!complete bra.uni waitLoop; | |
| } | |
| // end inline asm | |
| .loc 1 188 34 // bar.py:188:34 | |
| cp.async.wait_group 0; | |
| bar.sync 0; | |
| .loc 1 189 41 // bar.py:189:41 | |
| and.b32 %r566, %r662, 6291456; | |
| add.s32 %r564, %r566, %r2; | |
| .loc 1 188 34 // bar.py:188:34 | |
| ld.shared.v2.b64 {%rd462, %rd463}, [%r565+12288]; | |
| ld.shared.v2.b64 {%rd464, %rd465}, [%r565+12304]; | |
| ld.shared.v2.b64 {%rd466, %rd467}, [%r565+12320]; | |
| ld.shared.v2.b64 {%rd468, %rd469}, [%r565+12336]; | |
| ld.shared.v2.b64 {%rd470, %rd471}, [%r565+12352]; | |
| ld.shared.v2.b64 {%rd472, %rd473}, [%r565+12368]; | |
| ld.shared.v2.b64 {%rd474, %rd475}, [%r565+12384]; | |
| ld.shared.v2.b64 {%rd476, %rd477}, [%r565+12400]; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| tcgen05.ld.sync.aligned.16x32bx2.x32.b32 {%r532, %r533, %r534, %r535, %r536, %r537, %r538, %r539, %r540, %r541, %r542, %r543, %r544, %r545, %r546, %r547, %r548, %r549, %r550, %r551, %r552, %r553, %r554, %r555, %r556, %r557, %r558, %r559, %r560, %r561, %r562, %r563}, [%r564 + 0], 32; | |
| // end inline asm | |
| cvt.u64.u32 %rd478, %r562; | |
| cvt.u64.u32 %rd479, %r563; | |
| shl.b64 %rd480, %rd479, 32; | |
| or.b64 %rd481, %rd478, %rd480; | |
| cvt.u64.u32 %rd482, %r560; | |
| cvt.u64.u32 %rd483, %r561; | |
| shl.b64 %rd484, %rd483, 32; | |
| or.b64 %rd485, %rd482, %rd484; | |
| cvt.u64.u32 %rd486, %r558; | |
| cvt.u64.u32 %rd487, %r559; | |
| shl.b64 %rd488, %rd487, 32; | |
| or.b64 %rd489, %rd486, %rd488; | |
| cvt.u64.u32 %rd490, %r556; | |
| cvt.u64.u32 %rd491, %r557; | |
| shl.b64 %rd492, %rd491, 32; | |
| or.b64 %rd493, %rd490, %rd492; | |
| cvt.u64.u32 %rd494, %r554; | |
| cvt.u64.u32 %rd495, %r555; | |
| shl.b64 %rd496, %rd495, 32; | |
| or.b64 %rd497, %rd494, %rd496; | |
| cvt.u64.u32 %rd498, %r552; | |
| cvt.u64.u32 %rd499, %r553; | |
| shl.b64 %rd500, %rd499, 32; | |
| or.b64 %rd501, %rd498, %rd500; | |
| cvt.u64.u32 %rd502, %r550; | |
| cvt.u64.u32 %rd503, %r551; | |
| shl.b64 %rd504, %rd503, 32; | |
| or.b64 %rd505, %rd502, %rd504; | |
| cvt.u64.u32 %rd506, %r548; | |
| cvt.u64.u32 %rd507, %r549; | |
| shl.b64 %rd508, %rd507, 32; | |
| or.b64 %rd509, %rd506, %rd508; | |
| cvt.u64.u32 %rd510, %r546; | |
| cvt.u64.u32 %rd511, %r547; | |
| shl.b64 %rd512, %rd511, 32; | |
| or.b64 %rd513, %rd510, %rd512; | |
| cvt.u64.u32 %rd514, %r544; | |
| cvt.u64.u32 %rd515, %r545; | |
| shl.b64 %rd516, %rd515, 32; | |
| or.b64 %rd517, %rd514, %rd516; | |
| cvt.u64.u32 %rd518, %r542; | |
| cvt.u64.u32 %rd519, %r543; | |
| shl.b64 %rd520, %rd519, 32; | |
| or.b64 %rd521, %rd518, %rd520; | |
| cvt.u64.u32 %rd522, %r540; | |
| cvt.u64.u32 %rd523, %r541; | |
| shl.b64 %rd524, %rd523, 32; | |
| or.b64 %rd525, %rd522, %rd524; | |
| cvt.u64.u32 %rd526, %r538; | |
| cvt.u64.u32 %rd527, %r539; | |
| shl.b64 %rd528, %rd527, 32; | |
| or.b64 %rd529, %rd526, %rd528; | |
| cvt.u64.u32 %rd530, %r536; | |
| cvt.u64.u32 %rd531, %r537; | |
| shl.b64 %rd532, %rd531, 32; | |
| or.b64 %rd533, %rd530, %rd532; | |
| cvt.u64.u32 %rd534, %r534; | |
| cvt.u64.u32 %rd535, %r535; | |
| shl.b64 %rd536, %rd535, 32; | |
| or.b64 %rd537, %rd534, %rd536; | |
| cvt.u64.u32 %rd538, %r532; | |
| cvt.u64.u32 %rd539, %r533; | |
| shl.b64 %rd540, %rd539, 32; | |
| or.b64 %rd541, %rd538, %rd540; | |
| tcgen05.wait::ld.sync.aligned; | |
| .loc 1 189 46 // bar.py:189:46 | |
| cvt.u64.u32 %rd542, %r498; | |
| shl.b64 %rd543, %rd542, 32; | |
| or.b64 %rd544, %rd542, %rd543; | |
| mul.f32x2 %rd545, %rd544, %rd541; | |
| mul.f32x2 %rd546, %rd544, %rd537; | |
| mul.f32x2 %rd547, %rd544, %rd533; | |
| mul.f32x2 %rd548, %rd544, %rd529; | |
| mul.f32x2 %rd549, %rd544, %rd525; | |
| mul.f32x2 %rd550, %rd544, %rd521; | |
| mul.f32x2 %rd551, %rd544, %rd517; | |
| mul.f32x2 %rd552, %rd544, %rd513; | |
| mul.f32x2 %rd553, %rd544, %rd509; | |
| mul.f32x2 %rd554, %rd544, %rd505; | |
| mul.f32x2 %rd555, %rd544, %rd501; | |
| mul.f32x2 %rd556, %rd544, %rd497; | |
| mul.f32x2 %rd557, %rd544, %rd493; | |
| mul.f32x2 %rd558, %rd544, %rd489; | |
| mul.f32x2 %rd559, %rd544, %rd485; | |
| mul.f32x2 %rd560, %rd544, %rd481; | |
| .loc 1 189 31 // bar.py:189:31 | |
| fma.rn.f32x2 %rd647, %rd477, %rd560, %rd647; | |
| fma.rn.f32x2 %rd646, %rd476, %rd559, %rd646; | |
| fma.rn.f32x2 %rd645, %rd475, %rd558, %rd645; | |
| fma.rn.f32x2 %rd644, %rd474, %rd557, %rd644; | |
| fma.rn.f32x2 %rd643, %rd473, %rd556, %rd643; | |
| fma.rn.f32x2 %rd642, %rd472, %rd555, %rd642; | |
| fma.rn.f32x2 %rd641, %rd471, %rd554, %rd641; | |
| fma.rn.f32x2 %rd640, %rd470, %rd553, %rd640; | |
| fma.rn.f32x2 %rd639, %rd469, %rd552, %rd639; | |
| fma.rn.f32x2 %rd638, %rd468, %rd551, %rd638; | |
| fma.rn.f32x2 %rd637, %rd467, %rd550, %rd637; | |
| fma.rn.f32x2 %rd636, %rd466, %rd549, %rd636; | |
| fma.rn.f32x2 %rd635, %rd465, %rd548, %rd635; | |
| fma.rn.f32x2 %rd634, %rd464, %rd547, %rd634; | |
| fma.rn.f32x2 %rd633, %rd463, %rd546, %rd633; | |
| fma.rn.f32x2 %rd632, %rd462, %rd545, %rd632; | |
| $L__BB0_114: | |
| .loc 1 0 31 // bar.py:0:31 | |
| cvt.u32.u64 %r584, %rd37; | |
| cvt.u32.u64 %r585, %rd36; | |
| cvt.u32.u64 %r586, %rd35; | |
| cvt.u32.u64 %r587, %rd34; | |
| .loc 1 114 30 // bar.py:114:30 | |
| setp.lt.s32 %p65, %r584, %r21; | |
| setp.lt.s32 %p66, %r585, %r21; | |
| setp.lt.s32 %p67, %r586, %r21; | |
| setp.lt.s32 %p68, %r587, %r21; | |
| .loc 1 172 22 // bar.py:172:22 | |
| cp.async.wait_group 0; | |
| bar.sync 0; | |
| add.s32 %r567, %r29, 12544; | |
| // begin inline asm | |
| @%p60 mbarrier.inval.shared::cta.b64 [%r567]; | |
| // end inline asm | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p60 mbarrier.inval.shared::cta.b64 [%r194]; | |
| // end inline asm | |
| .loc 1 212 33 // bar.py:212:33 | |
| mov.b64 {%r588, %r589}, %rd632; | |
| cvt.rn.bf16x2.f32 %r590, %r589, %r588; | |
| mov.b64 {%r591, %r592}, %rd633; | |
| cvt.rn.bf16x2.f32 %r593, %r592, %r591; | |
| mov.b64 {%r594, %r595}, %rd634; | |
| cvt.rn.bf16x2.f32 %r596, %r595, %r594; | |
| mov.b64 {%r597, %r598}, %rd635; | |
| cvt.rn.bf16x2.f32 %r599, %r598, %r597; | |
| mov.b64 {%r600, %r601}, %rd636; | |
| cvt.rn.bf16x2.f32 %r602, %r601, %r600; | |
| mov.b64 {%r603, %r604}, %rd637; | |
| cvt.rn.bf16x2.f32 %r605, %r604, %r603; | |
| mov.b64 {%r606, %r607}, %rd638; | |
| cvt.rn.bf16x2.f32 %r608, %r607, %r606; | |
| mov.b64 {%r609, %r610}, %rd639; | |
| cvt.rn.bf16x2.f32 %r611, %r610, %r609; | |
| mov.b64 {%r612, %r613}, %rd640; | |
| cvt.rn.bf16x2.f32 %r614, %r613, %r612; | |
| mov.b64 {%r615, %r616}, %rd641; | |
| cvt.rn.bf16x2.f32 %r617, %r616, %r615; | |
| mov.b64 {%r618, %r619}, %rd642; | |
| cvt.rn.bf16x2.f32 %r620, %r619, %r618; | |
| mov.b64 {%r621, %r622}, %rd643; | |
| cvt.rn.bf16x2.f32 %r623, %r622, %r621; | |
| mov.b64 {%r624, %r625}, %rd644; | |
| cvt.rn.bf16x2.f32 %r626, %r625, %r624; | |
| mov.b64 {%r627, %r628}, %rd645; | |
| cvt.rn.bf16x2.f32 %r629, %r628, %r627; | |
| mov.b64 {%r630, %r631}, %rd646; | |
| cvt.rn.bf16x2.f32 %r632, %r631, %r630; | |
| mov.b64 {%r633, %r634}, %rd647; | |
| cvt.rn.bf16x2.f32 %r635, %r634, %r633; | |
| .loc 1 213 37 // bar.py:213:37 | |
| or.b32 %r636, %r663, %r8; | |
| .loc 1 214 33 // bar.py:214:33 | |
| cvt.s64.s32 %rd565, %r25; | |
| mul.lo.s64 %rd566, %rd34, %rd565; | |
| mul.lo.s64 %rd567, %rd35, %rd565; | |
| mul.lo.s64 %rd568, %rd36, %rd565; | |
| mul.lo.s64 %rd569, %rd37, %rd565; | |
| .loc 1 214 21 // bar.py:214:21 | |
| shl.b64 %rd570, %rd566, 1; | |
| add.s64 %rd571, %rd115, %rd570; | |
| shl.b64 %rd572, %rd567, 1; | |
| add.s64 %rd573, %rd115, %rd572; | |
| shl.b64 %rd574, %rd568, 1; | |
| add.s64 %rd575, %rd115, %rd574; | |
| shl.b64 %rd576, %rd569, 1; | |
| add.s64 %rd577, %rd115, %rd576; | |
| .loc 1 214 55 // bar.py:214:55 | |
| mul.wide.s32 %rd578, %r636, 2; | |
| add.s64 %rd561, %rd571, %rd578; | |
| add.s64 %rd562, %rd573, %rd578; | |
| add.s64 %rd563, %rd575, %rd578; | |
| add.s64 %rd564, %rd577, %rd578; | |
| .loc 1 215 55 // bar.py:215:55 | |
| setp.lt.s32 %p69, %r636, %r19; | |
| .loc 1 215 36 // bar.py:215:36 | |
| and.pred %p61, %p69, %p68; | |
| and.pred %p62, %p69, %p67; | |
| and.pred %p63, %p69, %p66; | |
| and.pred %p64, %p69, %p65; | |
| .loc 1 216 21 // bar.py:216:21 | |
| and.b32 %r637, %r15, 176; | |
| shl.b32 %r638, %r6, 3; | |
| bfe.s32 %r639, %r1, 2, 1; | |
| and.b32 %r640, %r639, 1088; | |
| shl.b32 %r641, %r9, 2; | |
| xor.b32 %r642, %r640, %r641; | |
| add.s32 %r643, %r29, %r637; | |
| add.s32 %r644, %r643, %r638; | |
| add.s32 %r645, %r644, %r642; | |
| st.shared.v4.b32 [%r645], {%r590, %r602, %r614, %r626}; | |
| bar.sync 0; | |
| shl.b32 %r646, %r1, 5; | |
| and.b32 %r647, %r646, 1792; | |
| and.b32 %r648, %r7, 48; | |
| shl.b32 %r649, %r6, 1; | |
| shl.b32 %r650, %r12, 6; | |
| xor.b32 %r651, %r649, %r650; | |
| add.s32 %r652, %r29, %r647; | |
| add.s32 %r653, %r652, %r648; | |
| add.s32 %r654, %r653, %r651; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r568, %r572, %r576, %r580}, [%r654]; | |
| bar.sync 0; | |
| st.shared.v4.b32 [%r645], {%r593, %r605, %r617, %r629}; | |
| bar.sync 0; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r569, %r573, %r577, %r581}, [%r654]; | |
| bar.sync 0; | |
| st.shared.v4.b32 [%r645], {%r596, %r608, %r620, %r632}; | |
| bar.sync 0; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r570, %r574, %r578, %r582}, [%r654]; | |
| bar.sync 0; | |
| st.shared.v4.b32 [%r645], {%r599, %r611, %r623, %r635}; | |
| bar.sync 0; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r571, %r575, %r579, %r583}, [%r654]; | |
| // begin inline asm | |
| @%p61 st.global.v4.b32 [ %rd561 + 0 ], { %r568, %r569, %r570, %r571 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p62 st.global.v4.b32 [ %rd562 + 0 ], { %r572, %r573, %r574, %r575 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p63 st.global.v4.b32 [ %rd563 + 0 ], { %r576, %r577, %r578, %r579 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p64 st.global.v4.b32 [ %rd564 + 0 ], { %r580, %r581, %r582, %r583 }; | |
| // end inline asm | |
| .loc 1 216 4 // bar.py:216:4 | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r2, 64; | |
| // end inline asm | |
| $L__BB0_2: // %common.ret | |
| .loc 1 0 0 // bar.py:0 | |
| ret; | |
| $L__tmp8: | |
| $L__func_end0: | |
| // -- End function | |
| } | |
| .file 1 "/home/dev/bar.py" | |
| .file 2 "/home/dev/py3.12-torch2.12/lib/python3.12/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 0 // EOM(3) | |
| } | |
| .section .debug_info | |
| { | |
| .b32 176 // 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:0xa9 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 98 // DW_AT_name | |
| .b8 97 | |
| .b8 114 | |
| .b8 46 | |
| .b8 112 | |
| .b8 121 | |
| .b8 0 | |
| .b32 .debug_line // DW_AT_stmt_list | |
| .b8 47 // DW_AT_comp_dir | |
| .b8 104 | |
| .b8 111 | |
| .b8 109 | |
| .b8 101 | |
| .b8 47 | |
| .b8 100 | |
| .b8 101 | |
| .b8 118 | |
| .b8 0 | |
| .b8 2 // Abbrev [2] 0x2a:0x13 DW_TAG_subprogram | |
| .b8 102 // DW_AT_name | |
| .b8 117 | |
| .b8 115 | |
| .b8 101 | |
| .b8 100 | |
| .b8 95 | |
| .b8 109 | |
| .b8 111 | |
| .b8 101 | |
| .b8 95 | |
| .b8 107 | |
| .b8 101 | |
| .b8 114 | |
| .b8 110 | |
| .b8 101 | |
| .b8 108 | |
| .b8 0 | |
| .b8 1 // DW_AT_inline | |
| .b8 3 // Abbrev [3] 0x3d:0x76 DW_TAG_subprogram | |
| .b64 $L__func_begin0 // DW_AT_low_pc | |
| .b64 $L__func_end0 // DW_AT_high_pc | |
| .b32 42 // DW_AT_abstract_origin | |
| .b8 4 // Abbrev [4] 0x52:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // 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 94 // DW_AT_call_line | |
| .b8 28 // DW_AT_call_column | |
| .b8 4 // Abbrev [4] 0x6a:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // DW_AT_abstract_origin | |
| .b64 $L__tmp2 // DW_AT_low_pc | |
| .b64 $L__tmp3 // DW_AT_high_pc | |
| .b8 1 // DW_AT_call_file | |
| .b8 95 // DW_AT_call_line | |
| .b8 27 // DW_AT_call_column | |
| .b8 4 // Abbrev [4] 0x82:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // DW_AT_abstract_origin | |
| .b64 $L__tmp4 // DW_AT_low_pc | |
| .b64 $L__tmp5 // DW_AT_high_pc | |
| .b8 1 // DW_AT_call_file | |
| .b8 128 // DW_AT_call_line | |
| .b8 12 // DW_AT_call_column | |
| .b8 4 // Abbrev [4] 0x9a:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // DW_AT_abstract_origin | |
| .b64 $L__tmp6 // DW_AT_low_pc | |
| .b64 $L__tmp7 // DW_AT_high_pc | |
| .b8 1 // DW_AT_call_file | |
| .b8 172 // DW_AT_call_line | |
| .b8 33 // DW_AT_call_column | |
| .b8 0 // End Of Children Mark | |
| .b8 0 // End Of Children Mark | |
| } | |
| .section .debug_macinfo { } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // | |
| // Generated by LLVM NVPTX Back-End | |
| // | |
| .version 9.1 | |
| .target sm_100a | |
| .address_size 64 | |
| // .globl fused_moe_kernel // -- Begin function fused_moe_kernel | |
| .extern .shared .align 16 .b8 global_smem[]; | |
| // @fused_moe_kernel | |
| .visible .entry fused_moe_kernel( | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_0, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_1, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_2, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_3, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_4, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_5, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_6, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_7, | |
| .param .u32 fused_moe_kernel_param_8, | |
| .param .u32 fused_moe_kernel_param_9, | |
| .param .u32 fused_moe_kernel_param_10, | |
| .param .u32 fused_moe_kernel_param_11, | |
| .param .u32 fused_moe_kernel_param_12, | |
| .param .u32 fused_moe_kernel_param_13, | |
| .param .u32 fused_moe_kernel_param_14, | |
| .param .u32 fused_moe_kernel_param_15, | |
| .param .u32 fused_moe_kernel_param_16, | |
| .param .u32 fused_moe_kernel_param_17, | |
| .param .u32 fused_moe_kernel_param_18, | |
| .param .u32 fused_moe_kernel_param_19, | |
| .param .u32 fused_moe_kernel_param_20, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_21, | |
| .param .u64 .ptr .global .align 1 fused_moe_kernel_param_22 | |
| ) | |
| .reqntid 128 | |
| { | |
| .reg .pred %p<57>; | |
| .reg .b32 %r<342>; | |
| .reg .b64 %rd<236>; | |
| .loc 1 47 0 // bar.py:47:0 | |
| $L__func_begin0: | |
| .loc 1 47 0 // bar.py:47:0 | |
| // %bb.0: | |
| ld.param.b32 %r19, [fused_moe_kernel_param_8]; | |
| $L__tmp0: | |
| .loc 1 47 0 // bar.py:47 | |
| mov.u32 %r1, %tid.x; | |
| setp.lt.u32 %p1, %r1, 32; | |
| mov.b32 %r29, global_smem; | |
| // begin inline asm | |
| @%p1 tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r29], 64; | |
| // end inline asm | |
| bar.sync 0; | |
| ld.shared.b32 %r2, [global_smem]; | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned; | |
| // end inline asm | |
| ld.param.b64 %rd27, [fused_moe_kernel_param_7]; | |
| .loc 1 93 24 // bar.py:93:24 | |
| mov.u32 %r31, %ctaid.x; | |
| ld.param.b32 %r32, [fused_moe_kernel_param_10]; | |
| $L__tmp1: | |
| .loc 2 43 17 // standard.py:43:17 @[ bar.py:94:28 ] | |
| add.s32 %r33, %r32, 63; | |
| .loc 2 43 30 // standard.py:43:30 @[ bar.py:94:28 ] | |
| shr.s32 %r34, %r33, 31; | |
| shr.u32 %r35, %r34, 26; | |
| add.s32 %r36, %r33, %r35; | |
| shr.s32 %r37, %r36, 6; | |
| $L__tmp2: | |
| .loc 2 43 17 // standard.py:43:17 @[ bar.py:95:27 ] | |
| add.s32 %r38, %r19, 63; | |
| .loc 2 43 30 // standard.py:43:30 @[ bar.py:95:27 ] | |
| shr.s32 %r39, %r38, 31; | |
| shr.u32 %r40, %r39, 26; | |
| add.s32 %r41, %r38, %r40; | |
| shr.s32 %r42, %r41, 6; | |
| $L__tmp3: | |
| .loc 1 96 38 // bar.py:96:38 | |
| shl.b32 %r43, %r42, 3; | |
| .loc 1 97 22 // bar.py:97:22 | |
| div.s32 %r44, %r31, %r43; | |
| .loc 1 98 29 // bar.py:98:29 | |
| shl.b32 %r45, %r44, 3; | |
| .loc 1 99 35 // bar.py:99:35 | |
| sub.s32 %r46, %r37, %r45; | |
| .loc 1 99 48 // bar.py:99:48 | |
| min.s32 %r47, %r46, 8; | |
| .loc 1 100 34 // bar.py:100:34 | |
| mul.lo.s32 %r48, %r44, %r43; | |
| sub.s32 %r49, %r31, %r48; | |
| .loc 1 101 40 // bar.py:101:40 | |
| div.s32 %r3, %r49, %r47; | |
| .loc 1 100 54 // bar.py:100:54 | |
| mul.lo.s32 %r50, %r3, %r47; | |
| sub.s32 %r51, %r49, %r50; | |
| .loc 1 100 27 // bar.py:100:27 | |
| add.s32 %r4, %r51, %r45; | |
| .loc 1 104 37 // bar.py:104:37 | |
| // begin inline asm | |
| mov.u32 %r30, 0x0; | |
| ld.global.b32 { %r30 }, [ %rd27 + 0 ]; | |
| // end inline asm | |
| .loc 1 105 15 // bar.py:105:15 | |
| shl.b32 %r9, %r4, 6; | |
| .loc 1 105 31 // bar.py:105:31 | |
| setp.lt.s32 %p2, %r9, %r30; | |
| @%p2 bra $L__BB0_3; | |
| bra.uni $L__BB0_1; | |
| $L__BB0_3: | |
| .loc 1 0 31 // bar.py:0:31 | |
| ld.param.b32 %r25, [fused_moe_kernel_param_15]; | |
| ld.param.b32 %r21, [fused_moe_kernel_param_11]; | |
| ld.param.b64 %rd26, [fused_moe_kernel_param_6]; | |
| ld.param.b64 %rd25, [fused_moe_kernel_param_5]; | |
| ld.param.b64 %rd22, [fused_moe_kernel_param_2]; | |
| .loc 1 103 0 // bar.py:103 | |
| bfe.u32 %r52, %r1, 1, 6; | |
| and.b32 %r6, %r1, 96; | |
| shl.b32 %r7, %r1, 3; | |
| and.b32 %r8, %r7, 56; | |
| cvt.u64.u32 %rd1, %r52; | |
| .loc 1 103 24 // bar.py:103:24 | |
| and.b32 %r61, %r1, 15; | |
| shr.u32 %r62, %r6, 1; | |
| or.b32 %r63, %r62, %r61; | |
| bfe.u32 %r64, %r1, 3, 4; | |
| .loc 1 109 52 // bar.py:109:52 | |
| mad.wide.s32 %rd35, %r9, 4, %rd25; | |
| mad.wide.u32 %rd28, %r64, 4, %rd35; | |
| add.s64 %rd29, %rd28, 64; | |
| add.s64 %rd30, %rd28, 128; | |
| add.s64 %rd31, %rd28, 192; | |
| shl.b64 %rd36, %rd1, 2; | |
| add.s64 %rd32, %rd35, %rd36; | |
| mad.wide.u32 %rd33, %r63, 4, %rd35; | |
| .loc 1 109 29 // bar.py:109:29 | |
| // begin inline asm | |
| mov.u32 %r54, 0x0; | |
| ld.global.b32 { %r54 }, [ %rd28 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r55, 0x0; | |
| ld.global.b32 { %r55 }, [ %rd29 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r56, 0x0; | |
| ld.global.b32 { %r56 }, [ %rd30 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r57, 0x0; | |
| ld.global.b32 { %r57 }, [ %rd31 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r58, 0x0; | |
| ld.global.b32 { %r58 }, [ %rd32 + 0 ]; | |
| // end inline asm | |
| // begin inline asm | |
| mov.u32 %r59, 0x0; | |
| ld.global.b32 { %r59 }, [ %rd33 + 0 ]; | |
| // end inline asm | |
| .loc 1 112 31 // bar.py:112:31 | |
| cvt.s64.s32 %rd3, %r54; | |
| cvt.s64.s32 %rd4, %r55; | |
| cvt.s64.s32 %rd5, %r56; | |
| cvt.s64.s32 %rd6, %r57; | |
| .loc 1 116 43 // bar.py:116:43 | |
| mad.wide.s32 %rd34, %r4, 4, %rd26; | |
| .loc 1 116 26 // bar.py:116:26 | |
| // begin inline asm | |
| mov.u32 %r60, 0x0; | |
| ld.global.b32 { %r60 }, [ %rd34 + 0 ]; | |
| // end inline asm | |
| .loc 1 117 22 // bar.py:117:22 | |
| setp.ne.b32 %p3, %r60, -1; | |
| @%p3 bra $L__BB0_5; | |
| bra.uni $L__BB0_4; | |
| $L__BB0_5: | |
| .loc 1 0 0 // bar.py:0 | |
| and.b32 %r53, %r1, 63; | |
| cvt.u64.u32 %rd2, %r53; | |
| .loc 1 132 23 // bar.py:132:23 | |
| shl.b32 %r10, %r3, 6; | |
| .loc 1 132 38 // bar.py:132:38 | |
| cvt.s64.s32 %rd37, %r10; | |
| or.b64 %rd8, %rd37, %rd1; | |
| or.b64 %rd9, %rd37, %rd2; | |
| .loc 1 132 81 // bar.py:132:81 | |
| cvt.s64.s32 %rd10, %r19; | |
| or.b64 %rd38, %rd37, %rd10; | |
| and.b64 %rd39, %rd38, -4294967296; | |
| setp.ne.b64 %p4, %rd39, 0; | |
| @%p4 bra $L__BB0_7; | |
| bra.uni $L__BB0_6; | |
| $L__BB0_7: | |
| rem.s64 %rd216, %rd8, %rd10; | |
| bra.uni $L__BB0_8; | |
| $L__BB0_1: | |
| .loc 1 106 8 // bar.py:106:8 | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r2, 64; | |
| // end inline asm | |
| bra.uni $L__BB0_2; | |
| $L__BB0_4: | |
| .loc 1 0 8 // bar.py:0:8 | |
| cvt.u32.u64 %r323, %rd6; | |
| cvt.u32.u64 %r324, %rd5; | |
| cvt.u32.u64 %r325, %rd4; | |
| cvt.u32.u64 %r326, %rd3; | |
| .loc 1 114 30 // bar.py:114:30 | |
| setp.lt.s32 %p51, %r323, %r21; | |
| setp.lt.s32 %p52, %r324, %r21; | |
| setp.lt.s32 %p53, %r325, %r21; | |
| setp.lt.s32 %p54, %r326, %r21; | |
| $L__tmp4: | |
| .loc 1 40 22 // bar.py:40:22 @[ bar.py:128:12 ] | |
| shl.b32 %r327, %r3, 6; | |
| .loc 1 40 37 // bar.py:40:37 @[ bar.py:128:12 ] | |
| or.b32 %r328, %r327, %r8; | |
| .loc 1 41 33 // bar.py:41:33 @[ bar.py:128:12 ] | |
| cvt.s64.s32 %rd202, %r25; | |
| mul.lo.s64 %rd203, %rd3, %rd202; | |
| mul.lo.s64 %rd204, %rd4, %rd202; | |
| mul.lo.s64 %rd205, %rd5, %rd202; | |
| mul.lo.s64 %rd206, %rd6, %rd202; | |
| .loc 1 41 21 // bar.py:41:21 @[ bar.py:128:12 ] | |
| shl.b64 %rd207, %rd203, 1; | |
| add.s64 %rd208, %rd22, %rd207; | |
| shl.b64 %rd209, %rd204, 1; | |
| add.s64 %rd210, %rd22, %rd209; | |
| shl.b64 %rd211, %rd205, 1; | |
| add.s64 %rd212, %rd22, %rd211; | |
| shl.b64 %rd213, %rd206, 1; | |
| add.s64 %rd214, %rd22, %rd213; | |
| .loc 1 41 55 // bar.py:41:55 @[ bar.py:128:12 ] | |
| mul.wide.s32 %rd215, %r328, 2; | |
| add.s64 %rd198, %rd208, %rd215; | |
| add.s64 %rd199, %rd210, %rd215; | |
| add.s64 %rd200, %rd212, %rd215; | |
| add.s64 %rd201, %rd214, %rd215; | |
| .loc 1 42 55 // bar.py:42:55 @[ bar.py:128:12 ] | |
| setp.lt.s32 %p55, %r328, %r19; | |
| .loc 1 42 36 // bar.py:42:36 @[ bar.py:128:12 ] | |
| and.pred %p47, %p55, %p54; | |
| and.pred %p48, %p55, %p53; | |
| and.pred %p49, %p55, %p52; | |
| and.pred %p50, %p55, %p51; | |
| mov.b32 %r322, 0; | |
| .loc 1 43 21 // bar.py:43:21 @[ bar.py:128:12 ] | |
| // begin inline asm | |
| @%p47 st.global.v4.b32 [ %rd198 + 0 ], { %r322, %r322, %r322, %r322 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p48 st.global.v4.b32 [ %rd199 + 0 ], { %r322, %r322, %r322, %r322 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p49 st.global.v4.b32 [ %rd200 + 0 ], { %r322, %r322, %r322, %r322 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p50 st.global.v4.b32 [ %rd201 + 0 ], { %r322, %r322, %r322, %r322 }; | |
| // end inline asm | |
| $L__tmp5: | |
| .loc 1 130 8 // bar.py:130:8 | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r2, 64; | |
| // end inline asm | |
| bra.uni $L__BB0_2; | |
| $L__BB0_6: | |
| .loc 1 132 81 // bar.py:132:81 | |
| cvt.u32.u64 %r65, %rd10; | |
| cvt.u32.u64 %r66, %rd8; | |
| rem.u32 %r67, %r66, %r65; | |
| cvt.u64.u32 %rd216, %r67; | |
| $L__BB0_8: | |
| .loc 1 0 81 // bar.py:0:81 | |
| ld.param.b32 %r28, [fused_moe_kernel_param_18]; | |
| ld.param.b32 %r27, [fused_moe_kernel_param_17]; | |
| ld.param.b32 %r26, [fused_moe_kernel_param_16]; | |
| ld.param.b32 %r24, [fused_moe_kernel_param_14]; | |
| ld.param.b32 %r23, [fused_moe_kernel_param_13]; | |
| ld.param.b32 %r22, [fused_moe_kernel_param_12]; | |
| ld.param.b32 %r20, [fused_moe_kernel_param_9]; | |
| ld.param.b64 %rd24, [fused_moe_kernel_param_4]; | |
| ld.param.b64 %rd21, [fused_moe_kernel_param_1]; | |
| ld.param.b64 %rd20, [fused_moe_kernel_param_0]; | |
| shr.u32 %r5, %r1, 5; | |
| cvt.s64.s32 %rd7, %r60; | |
| .loc 1 132 81 // bar.py:132:81 | |
| or.b64 %rd40, %rd9, %rd10; | |
| and.b64 %rd41, %rd40, -4294967296; | |
| setp.ne.b64 %p5, %rd41, 0; | |
| @%p5 bra $L__BB0_10; | |
| bra.uni $L__BB0_9; | |
| $L__BB0_10: | |
| rem.s64 %rd217, %rd9, %rd10; | |
| bra.uni $L__BB0_11; | |
| $L__BB0_9: | |
| cvt.u32.u64 %r68, %rd10; | |
| cvt.u32.u64 %r69, %rd9; | |
| rem.u32 %r70, %r69, %r68; | |
| cvt.u64.u32 %rd217, %r70; | |
| $L__BB0_11: | |
| .loc 1 0 81 // bar.py:0:81 | |
| ld.param.b64 %rd23, [fused_moe_kernel_param_3]; | |
| .loc 1 114 30 // bar.py:114:30 | |
| setp.lt.s32 %p8, %r58, %r21; | |
| .loc 1 135 31 // bar.py:135:31 | |
| shr.s32 %r80, %r58, 31; | |
| shr.u32 %r81, %r80, 29; | |
| add.s32 %r82, %r58, %r81; | |
| shr.s32 %r83, %r82, 3; | |
| .loc 1 135 39 // bar.py:135:39 | |
| mul.wide.s32 %rd11, %r83, %r22; | |
| .loc 1 135 58 // bar.py:135:58 | |
| and.b32 %r11, %r1, 1; | |
| shl.b32 %r84, %r11, 4; | |
| .loc 1 135 51 // bar.py:135:51 | |
| cvt.u64.u32 %rd12, %r84; | |
| .loc 1 135 8 // bar.py:135:8 | |
| add.s64 %rd46, %rd20, %rd11; | |
| add.s64 %rd42, %rd46, %rd12; | |
| .loc 1 139 24 // bar.py:139:24 | |
| cvt.s64.s32 %rd47, %r23; | |
| mul.lo.s64 %rd13, %rd7, %rd47; | |
| .loc 1 139 10 // bar.py:139:10 | |
| add.s64 %rd48, %rd21, %rd13; | |
| .loc 1 140 60 // bar.py:140:60 | |
| cvt.s64.s32 %rd49, %r24; | |
| mul.lo.s64 %rd14, %rd216, %rd49; | |
| .loc 1 140 11 // bar.py:140:11 | |
| add.s64 %rd50, %rd48, %rd14; | |
| add.s64 %rd43, %rd50, %rd12; | |
| .loc 1 151 56 // bar.py:151:56 | |
| shr.s32 %r85, %r59, 31; | |
| shr.u32 %r86, %r85, 29; | |
| add.s32 %r87, %r59, %r86; | |
| shr.s32 %r88, %r87, 3; | |
| .loc 1 151 65 // bar.py:151:65 | |
| mul.wide.s32 %rd51, %r88, %r26; | |
| .loc 1 151 41 // bar.py:151:41 | |
| shl.b64 %rd52, %rd51, 2; | |
| .loc 1 152 34 // bar.py:152:34 | |
| cvt.u32.u64 %r89, %rd217; | |
| shr.s32 %r90, %r89, 31; | |
| shr.u32 %r91, %r90, 25; | |
| add.s32 %r92, %r89, %r91; | |
| shr.s32 %r93, %r92, 7; | |
| .loc 1 154 44 // bar.py:154:44 | |
| cvt.s64.s32 %rd53, %r27; | |
| mul.lo.s64 %rd54, %rd7, %rd53; | |
| .loc 1 154 30 // bar.py:154:30 | |
| shl.b64 %rd55, %rd54, 2; | |
| add.s64 %rd56, %rd24, %rd55; | |
| .loc 1 154 68 // bar.py:154:68 | |
| mul.wide.s32 %rd57, %r93, %r28; | |
| .loc 1 154 57 // bar.py:154:57 | |
| shl.b64 %rd58, %rd57, 2; | |
| $L__tmp6: | |
| .loc 2 43 17 // standard.py:43:17 @[ bar.py:172:33 ] | |
| add.s32 %r12, %r20, 31; | |
| $L__tmp7: | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.eq.b32 %p9, %r5, 0; | |
| elect.sync %r97|%p10, -1; | |
| and.pred %p6, %p9, %p10; | |
| add.s32 %r340, %r29, 12544; | |
| // begin inline asm | |
| @%p6 mbarrier.init.shared::cta.b64 [%r340], 1; | |
| // end inline asm | |
| elect.sync %r98|%p11, -1; | |
| and.pred %p7, %p9, %p11; | |
| add.s32 %r71, %r29, 12552; | |
| // begin inline asm | |
| @%p7 mbarrier.init.shared::cta.b64 [%r71], 1; | |
| // end inline asm | |
| setp.lt.s32 %p12, %r12, 32; | |
| setp.gt.s32 %p13, %r12, 31; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p14, %r84, %r20; | |
| .loc 1 174 12 // bar.py:174:12 | |
| shl.b32 %r14, %r1, 4; | |
| and.b32 %r99, %r14, 1904; | |
| bfe.s32 %r100, %r1, 3, 1; | |
| and.b32 %r101, %r100, 144; | |
| xor.b32 %r102, %r101, %r99; | |
| add.s32 %r72, %r29, %r102; | |
| selp.b32 %r103, 16, 0, %p8; | |
| selp.b32 %r104, %r103, 0, %p14; | |
| selp.b32 %r73, %r104, 0, %p13; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r72 + 0 ], [ %rd42 + 0 ], 0x10, %r73; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r74, %r72, 6144; | |
| selp.b32 %r105, 16, 0, %p13; | |
| selp.b32 %r75, %r105, 0, %p14; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r74 + 0 ], [ %rd43 + 0 ], 0x10, %r75; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.gt.s32 %p15, %r12, 63; | |
| .loc 1 197 18 // bar.py:197:18 | |
| add.s64 %rd44, %rd42, 32; | |
| .loc 1 198 18 // bar.py:198:18 | |
| add.s64 %rd45, %rd43, 32; | |
| .loc 1 175 62 // bar.py:175:62 | |
| add.s32 %r106, %r20, -32; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p16, %r84, %r106; | |
| .loc 1 174 12 // bar.py:174:12 | |
| bar.sync 0; | |
| add.s32 %r76, %r72, 2048; | |
| selp.b32 %r107, %r103, 0, %p16; | |
| selp.b32 %r77, %r107, 0, %p15; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r76 + 0 ], [ %rd44 + 0 ], 0x10, %r77; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r78, %r72, 8192; | |
| selp.b32 %r108, 16, 0, %p16; | |
| selp.b32 %r79, %r108, 0, %p15; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r78 + 0 ], [ %rd45 + 0 ], 0x10, %r79; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 174 12 // bar.py:174:12 | |
| cp.async.wait_group 2; | |
| bar.sync 0; | |
| .loc 1 189 41 // bar.py:189:41 | |
| shfl.sync.idx.b32 %r15, %r5, 0, 31, -1; | |
| setp.ne.b32 %p17, %r15, 0; | |
| or.pred %p18, %p12, %p17; | |
| @%p18 bra $L__BB0_13; | |
| // %bb.12: | |
| elect.sync %r111|%p20, -1; | |
| bfe.u32 %r112, %r29, 4, 14; | |
| cvt.u64.u32 %rd61, %r112; | |
| add.s32 %r113, %r29, 6144; | |
| bfe.u32 %r114, %r113, 4, 14; | |
| cvt.u64.u32 %rd62, %r114; | |
| or.b64 %rd59, %rd61, -4611615580963733504; | |
| or.b64 %rd60, %rd62, -4611615580963733504; | |
| mov.b32 %r109, 68157456; | |
| mov.pred %p19, 0; | |
| // begin inline asm | |
| @%p20 tcgen05.mma.cta_group::1.kind::f8f6f4 [ %r2 + 0 ], %rd59, %rd60, %r109, %p19; | |
| // end inline asm | |
| add.s32 %r110, %r29, 12544; | |
| // begin inline asm | |
| @%p20 tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%r110]; | |
| // end inline asm | |
| $L__BB0_13: | |
| .loc 1 0 41 // bar.py:0:41 | |
| add.s64 %rd15, %rd23, %rd52; | |
| add.s64 %rd16, %rd56, %rd58; | |
| cvt.u32.u64 %r119, %rd12; | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.gt.s32 %p21, %r12, 95; | |
| setp.lt.s32 %p22, %r12, 64; | |
| .loc 1 197 18 // bar.py:197:18 | |
| add.s64 %rd63, %rd42, 64; | |
| .loc 1 198 18 // bar.py:198:18 | |
| add.s64 %rd64, %rd43, 64; | |
| .loc 1 175 62 // bar.py:175:62 | |
| add.s32 %r120, %r20, -64; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p23, %r119, %r120; | |
| .loc 1 174 12 // bar.py:174:12 | |
| bar.sync 0; | |
| add.s32 %r115, %r72, 4096; | |
| selp.b32 %r121, %r103, 0, %p23; | |
| selp.b32 %r116, %r121, 0, %p21; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r115 + 0 ], [ %rd63 + 0 ], 0x10, %r116; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r117, %r72, 10240; | |
| selp.b32 %r122, 16, 0, %p23; | |
| selp.b32 %r118, %r122, 0, %p21; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r117 + 0 ], [ %rd64 + 0 ], 0x10, %r118; | |
| // end inline asm | |
| cp.async.commit_group; | |
| mov.b32 %r123, 0f00000000; | |
| mov.b64 %rd220, {%r123, %r123}; | |
| mov.b32 %r341, 0; | |
| mov.b64 %rd219, 0; | |
| cvt.u32.u64 %r329, %rd2; | |
| and.b32 %r330, %r1, 64; | |
| and.b32 %r331, %r7, 128; | |
| shl.b32 %r332, %r15, 21; | |
| setp.lt.s32 %p56, %r59, %r21; | |
| mov.b64 %rd221, %rd220; | |
| mov.b64 %rd222, %rd220; | |
| mov.b64 %rd223, %rd220; | |
| mov.b64 %rd224, %rd220; | |
| mov.b64 %rd225, %rd220; | |
| mov.b64 %rd226, %rd220; | |
| mov.b64 %rd227, %rd220; | |
| mov.b64 %rd228, %rd220; | |
| mov.b64 %rd229, %rd220; | |
| mov.b64 %rd230, %rd220; | |
| mov.b64 %rd231, %rd220; | |
| mov.b64 %rd232, %rd220; | |
| mov.b64 %rd233, %rd220; | |
| mov.b64 %rd234, %rd220; | |
| mov.b64 %rd235, %rd220; | |
| .loc 1 172 22 // bar.py:172:22 | |
| @%p22 bra $L__BB0_19; | |
| // %bb.14: // %.lr.ph | |
| .loc 1 0 22 // bar.py:0:22 | |
| shr.s32 %r94, %r12, 31; | |
| shr.u32 %r95, %r94, 27; | |
| add.s32 %r96, %r12, %r95; | |
| shr.s32 %r13, %r96, 5; | |
| add.s32 %r16, %r13, -3; | |
| shl.b32 %r125, %r329, 2; | |
| and.b32 %r126, %r332, 6291456; | |
| add.s32 %r163, %r126, %r2; | |
| .loc 1 172 22 // bar.py:172:22 | |
| max.s32 %r127, %r13, 2; | |
| add.s32 %r128, %r29, %r125; | |
| add.s32 %r130, %r128, 12288; | |
| add.s32 %r17, %r29, %r331; | |
| add.s64 %rd65, %rd14, %rd13; | |
| add.s64 %rd66, %rd65, %rd12; | |
| add.s64 %rd67, %rd66, %rd21; | |
| add.s64 %rd17, %rd67, 96; | |
| add.s64 %rd68, %rd11, %rd12; | |
| add.s64 %rd69, %rd68, %rd20; | |
| add.s64 %rd18, %rd69, 96; | |
| mad.wide.u32 %rd19, %r127, 32, -32; | |
| mov.b64 %rd220, {%r123, %r123}; | |
| add.s32 %r340, %r29, 12544; | |
| add.s32 %r333, %r20, -96; | |
| mov.b32 %r337, 1; | |
| mov.b32 %r336, 2; | |
| mov.b32 %r124, 0; | |
| mov.b64 %rd218, 0; | |
| setp.eq.b32 %p24, %r330, 0; | |
| mov.b32 %r334, %r124; | |
| mov.b32 %r335, %r124; | |
| mov.b32 %r338, %r124; | |
| mov.b32 %r339, %r124; | |
| mov.b64 %rd221, %rd220; | |
| mov.b64 %rd222, %rd220; | |
| mov.b64 %rd223, %rd220; | |
| mov.b64 %rd224, %rd220; | |
| mov.b64 %rd225, %rd220; | |
| mov.b64 %rd226, %rd220; | |
| mov.b64 %rd227, %rd220; | |
| mov.b64 %rd228, %rd220; | |
| mov.b64 %rd229, %rd220; | |
| mov.b64 %rd230, %rd220; | |
| mov.b64 %rd231, %rd220; | |
| mov.b64 %rd232, %rd220; | |
| mov.b64 %rd233, %rd220; | |
| mov.b64 %rd234, %rd220; | |
| mov.b64 %rd235, %rd220; | |
| bra.uni $L__BB0_15; | |
| $L__BB0_17: // in Loop: Header=BB0_15 Depth=1 | |
| .loc 1 0 0 // bar.py:0 | |
| fma.rn.f32x2 %rd234, %rd87, %rd121, %rd234; | |
| fma.rn.f32x2 %rd233, %rd86, %rd120, %rd233; | |
| fma.rn.f32x2 %rd232, %rd85, %rd119, %rd232; | |
| fma.rn.f32x2 %rd231, %rd84, %rd118, %rd231; | |
| fma.rn.f32x2 %rd230, %rd83, %rd117, %rd230; | |
| fma.rn.f32x2 %rd229, %rd82, %rd116, %rd229; | |
| fma.rn.f32x2 %rd228, %rd81, %rd115, %rd228; | |
| fma.rn.f32x2 %rd227, %rd80, %rd114, %rd227; | |
| fma.rn.f32x2 %rd226, %rd79, %rd113, %rd226; | |
| fma.rn.f32x2 %rd225, %rd78, %rd112, %rd225; | |
| fma.rn.f32x2 %rd224, %rd77, %rd111, %rd224; | |
| fma.rn.f32x2 %rd223, %rd76, %rd110, %rd223; | |
| fma.rn.f32x2 %rd222, %rd75, %rd109, %rd222; | |
| fma.rn.f32x2 %rd221, %rd74, %rd108, %rd221; | |
| fma.rn.f32x2 %rd220, %rd73, %rd107, %rd220; | |
| fma.rn.f32x2 %rd235, %rd88, %rd106, %rd235; | |
| .loc 1 172 22 // bar.py:172:22 | |
| setp.lt.s32 %p28, %r339, %r16; | |
| .loc 1 189 41 // bar.py:189:41 | |
| add.s32 %r180, %r337, 1; | |
| setp.gt.s32 %p29, %r180, 1; | |
| selp.b32 %r337, 0, %r180, %p29; | |
| selp.b32 %r181, 1, 0, %p29; | |
| xor.b32 %r338, %r341, %r181; | |
| .loc 1 198 18 // bar.py:198:18 | |
| add.s64 %rd126, %rd18, %rd218; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s64 %rd127, %rd17, %rd218; | |
| add.s32 %r182, %r336, 1; | |
| setp.gt.s32 %p30, %r182, 2; | |
| selp.b32 %r336, 0, %r182, %p30; | |
| .loc 1 175 58 // bar.py:175:58 | |
| setp.lt.s32 %p31, %r119, %r333; | |
| .loc 1 174 12 // bar.py:174:12 | |
| shl.b32 %r183, %r336, 11; | |
| bar.sync 0; | |
| add.s32 %r176, %r72, %r183; | |
| selp.b32 %r184, 16, 0, %p31; | |
| selp.b32 %r185, %r184, 0, %p8; | |
| selp.b32 %r177, %r185, 0, %p28; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r176 + 0 ], [ %rd126 + 0 ], 0x10, %r177; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r178, %r176, 6144; | |
| selp.b32 %r179, %r184, 0, %p28; | |
| // begin inline asm | |
| cp.async.cg.shared.global [ %r178 + 0 ], [ %rd127 + 0 ], 0x10, %r179; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s32 %r339, %r339, 1; | |
| add.s64 %rd218, %rd218, 32; | |
| cvt.u32.u64 %r18, %rd218; | |
| cvt.u32.u64 %r186, %rd19; | |
| add.s32 %r333, %r333, -32; | |
| setp.ne.b32 %p32, %r186, %r18; | |
| mov.b32 %r334, %r341; | |
| @%p32 bra $L__BB0_15; | |
| bra.uni $L__BB0_18; | |
| $L__BB0_15: // =>This Inner Loop Header: Depth=1 | |
| .loc 1 0 22 // bar.py:0:22 | |
| mov.b32 %r341, %r338; | |
| cvt.u32.u64 %r164, %rd218; | |
| .loc 1 184 37 // bar.py:184:37 | |
| shr.u32 %r165, %r164, 7; | |
| .loc 1 186 35 // bar.py:186:35 | |
| mul.wide.u32 %rd72, %r165, 4; | |
| add.s64 %rd70, %rd15, %rd72; | |
| .loc 1 186 20 // bar.py:186:20 | |
| // begin inline asm | |
| mov.u32 %r129, %r124; | |
| @%p56 ld.global.b32 { %r129 }, [ %rd70 + 0 ]; | |
| // end inline asm | |
| .loc 1 188 49 // bar.py:188:49 | |
| add.s64 %rd71, %rd16, %rd72; | |
| .loc 1 188 34 // bar.py:188:34 | |
| // begin inline asm | |
| @%p24 cp.async.ca.shared.global [ %r130 + 0 ], [ %rd71 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| { | |
| .reg .pred complete; | |
| waitLoop: | |
| mbarrier.try_wait.parity.shared::cta.b64 complete, [%r340], %r334; | |
| @!complete bra.uni waitLoop; | |
| } | |
| // end inline asm | |
| .loc 1 188 34 // bar.py:188:34 | |
| cp.async.wait_group 0; | |
| bar.sync 0; | |
| .loc 1 172 22 // bar.py:172:22 | |
| add.s32 %r166, %r335, 1; | |
| setp.gt.s32 %p25, %r166, 2; | |
| selp.b32 %r335, 0, %r166, %p25; | |
| shl.b32 %r167, %r337, 3; | |
| add.s32 %r168, %r29, %r167; | |
| add.s32 %r340, %r168, 12544; | |
| .loc 1 188 34 // bar.py:188:34 | |
| ld.shared.v2.b64 {%rd73, %rd74}, [%r17+12288]; | |
| ld.shared.v2.b64 {%rd75, %rd76}, [%r17+12304]; | |
| ld.shared.v2.b64 {%rd77, %rd78}, [%r17+12320]; | |
| ld.shared.v2.b64 {%rd79, %rd80}, [%r17+12336]; | |
| ld.shared.v2.b64 {%rd81, %rd82}, [%r17+12352]; | |
| ld.shared.v2.b64 {%rd83, %rd84}, [%r17+12368]; | |
| ld.shared.v2.b64 {%rd85, %rd86}, [%r17+12384]; | |
| ld.shared.v2.b64 {%rd87, %rd88}, [%r17+12400]; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| tcgen05.ld.sync.aligned.16x32bx2.x32.b32 {%r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162}, [%r163 + 0], 32; | |
| // end inline asm | |
| mov.b64 %rd89, {%r159, %r160}; | |
| mov.b64 %rd90, {%r157, %r158}; | |
| mov.b64 %rd91, {%r155, %r156}; | |
| mov.b64 %rd92, {%r153, %r154}; | |
| mov.b64 %rd93, {%r151, %r152}; | |
| mov.b64 %rd94, {%r149, %r150}; | |
| mov.b64 %rd95, {%r147, %r148}; | |
| mov.b64 %rd96, {%r145, %r146}; | |
| mov.b64 %rd97, {%r143, %r144}; | |
| mov.b64 %rd98, {%r141, %r142}; | |
| mov.b64 %rd99, {%r139, %r140}; | |
| mov.b64 %rd100, {%r137, %r138}; | |
| mov.b64 %rd101, {%r135, %r136}; | |
| mov.b64 %rd102, {%r133, %r134}; | |
| mov.b64 %rd103, {%r131, %r132}; | |
| mov.b64 %rd104, {%r161, %r162}; | |
| tcgen05.wait::ld.sync.aligned; | |
| .loc 1 189 46 // bar.py:189:46 | |
| mov.b64 %rd105, {%r129, %r129}; | |
| mul.f32x2 %rd106, %rd105, %rd104; | |
| mul.f32x2 %rd107, %rd105, %rd103; | |
| mul.f32x2 %rd108, %rd105, %rd102; | |
| mul.f32x2 %rd109, %rd105, %rd101; | |
| mul.f32x2 %rd110, %rd105, %rd100; | |
| mul.f32x2 %rd111, %rd105, %rd99; | |
| mul.f32x2 %rd112, %rd105, %rd98; | |
| mul.f32x2 %rd113, %rd105, %rd97; | |
| mul.f32x2 %rd114, %rd105, %rd96; | |
| mul.f32x2 %rd115, %rd105, %rd95; | |
| mul.f32x2 %rd116, %rd105, %rd94; | |
| mul.f32x2 %rd117, %rd105, %rd93; | |
| mul.f32x2 %rd118, %rd105, %rd92; | |
| mul.f32x2 %rd119, %rd105, %rd91; | |
| mul.f32x2 %rd120, %rd105, %rd90; | |
| mul.f32x2 %rd121, %rd105, %rd89; | |
| .loc 1 189 41 // bar.py:189:41 | |
| @%p17 bra $L__BB0_17; | |
| // %bb.16: // in Loop: Header=BB0_15 Depth=1 | |
| .loc 1 174 12 // bar.py:174:12 | |
| shl.b32 %r170, %r335, 11; | |
| .loc 1 178 20 // bar.py:178:20 | |
| add.s32 %r171, %r29, %r170; | |
| add.s32 %r172, %r171, 6144; | |
| .loc 1 189 41 // bar.py:189:41 | |
| elect.sync %r173|%p27, -1; | |
| bfe.u32 %r174, %r171, 4, 14; | |
| cvt.u64.u32 %rd124, %r174; | |
| bfe.u32 %r175, %r172, 4, 14; | |
| cvt.u64.u32 %rd125, %r175; | |
| or.b64 %rd122, %rd124, -4611615580963733504; | |
| or.b64 %rd123, %rd125, -4611615580963733504; | |
| mov.b32 %r169, 68157456; | |
| mov.pred %p26, 0; | |
| // begin inline asm | |
| @%p27 tcgen05.mma.cta_group::1.kind::f8f6f4 [ %r2 + 0 ], %rd122, %rd123, %r169, %p26; | |
| // end inline asm | |
| // begin inline asm | |
| @%p27 tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%r340]; | |
| // end inline asm | |
| bra.uni $L__BB0_17; | |
| $L__BB0_18: // %._crit_edge.loopexit | |
| .loc 1 184 37 // bar.py:184:37 | |
| shr.s32 %r187, %r18, 31; | |
| shr.u32 %r188, %r187, 25; | |
| add.s32 %r189, %r18, %r188; | |
| shr.s32 %r190, %r189, 7; | |
| .loc 1 186 35 // bar.py:186:35 | |
| cvt.s64.s32 %rd219, %r190; | |
| $L__BB0_19: // %._crit_edge | |
| .loc 1 172 22 // bar.py:172:22 | |
| @%p12 bra $L__BB0_21; | |
| // %bb.20: | |
| .loc 1 186 35 // bar.py:186:35 | |
| shl.b64 %rd130, %rd219, 2; | |
| add.s64 %rd128, %rd15, %rd130; | |
| mov.b32 %r192, 0; | |
| .loc 1 186 20 // bar.py:186:20 | |
| // begin inline asm | |
| mov.u32 %r191, %r192; | |
| @%p56 ld.global.b32 { %r191 }, [ %rd128 + 0 ]; | |
| // end inline asm | |
| .loc 1 188 49 // bar.py:188:49 | |
| add.s64 %rd129, %rd16, %rd130; | |
| .loc 1 188 34 // bar.py:188:34 | |
| setp.eq.b32 %p33, %r330, 0; | |
| shl.b32 %r227, %r329, 2; | |
| add.s32 %r228, %r29, %r227; | |
| add.s32 %r193, %r228, 12288; | |
| // begin inline asm | |
| @%p33 cp.async.ca.shared.global [ %r193 + 0 ], [ %rd129 + 0 ], 0x4, 0x4; | |
| // end inline asm | |
| cp.async.commit_group; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| { | |
| .reg .pred complete; | |
| waitLoop: | |
| mbarrier.try_wait.parity.shared::cta.b64 complete, [%r340], %r341; | |
| @!complete bra.uni waitLoop; | |
| } | |
| // end inline asm | |
| .loc 1 188 34 // bar.py:188:34 | |
| cp.async.wait_group 0; | |
| bar.sync 0; | |
| add.s32 %r229, %r29, %r331; | |
| .loc 1 189 41 // bar.py:189:41 | |
| and.b32 %r230, %r332, 6291456; | |
| add.s32 %r226, %r230, %r2; | |
| .loc 1 188 34 // bar.py:188:34 | |
| ld.shared.v2.b64 {%rd131, %rd132}, [%r229+12400]; | |
| ld.shared.v2.b64 {%rd133, %rd134}, [%r229+12288]; | |
| ld.shared.v2.b64 {%rd135, %rd136}, [%r229+12304]; | |
| ld.shared.v2.b64 {%rd137, %rd138}, [%r229+12320]; | |
| ld.shared.v2.b64 {%rd139, %rd140}, [%r229+12336]; | |
| ld.shared.v2.b64 {%rd141, %rd142}, [%r229+12352]; | |
| ld.shared.v2.b64 {%rd143, %rd144}, [%r229+12368]; | |
| ld.shared.v2.b64 {%rd145, %rd146}, [%r229+12384]; | |
| .loc 1 189 41 // bar.py:189:41 | |
| // begin inline asm | |
| tcgen05.ld.sync.aligned.16x32bx2.x32.b32 {%r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225}, [%r226 + 0], 32; | |
| // end inline asm | |
| mov.b64 %rd147, {%r220, %r221}; | |
| mov.b64 %rd148, {%r218, %r219}; | |
| mov.b64 %rd149, {%r216, %r217}; | |
| mov.b64 %rd150, {%r214, %r215}; | |
| mov.b64 %rd151, {%r212, %r213}; | |
| mov.b64 %rd152, {%r210, %r211}; | |
| mov.b64 %rd153, {%r208, %r209}; | |
| mov.b64 %rd154, {%r206, %r207}; | |
| mov.b64 %rd155, {%r204, %r205}; | |
| mov.b64 %rd156, {%r202, %r203}; | |
| mov.b64 %rd157, {%r200, %r201}; | |
| mov.b64 %rd158, {%r198, %r199}; | |
| mov.b64 %rd159, {%r196, %r197}; | |
| mov.b64 %rd160, {%r194, %r195}; | |
| mov.b64 %rd161, {%r222, %r223}; | |
| mov.b64 %rd162, {%r224, %r225}; | |
| tcgen05.wait::ld.sync.aligned; | |
| .loc 1 189 46 // bar.py:189:46 | |
| mov.b64 %rd163, {%r191, %r191}; | |
| mul.f32x2 %rd164, %rd163, %rd162; | |
| mul.f32x2 %rd165, %rd163, %rd161; | |
| mul.f32x2 %rd166, %rd163, %rd160; | |
| mul.f32x2 %rd167, %rd163, %rd159; | |
| mul.f32x2 %rd168, %rd163, %rd158; | |
| mul.f32x2 %rd169, %rd163, %rd157; | |
| mul.f32x2 %rd170, %rd163, %rd156; | |
| mul.f32x2 %rd171, %rd163, %rd155; | |
| mul.f32x2 %rd172, %rd163, %rd154; | |
| mul.f32x2 %rd173, %rd163, %rd153; | |
| mul.f32x2 %rd174, %rd163, %rd152; | |
| mul.f32x2 %rd175, %rd163, %rd151; | |
| mul.f32x2 %rd176, %rd163, %rd150; | |
| mul.f32x2 %rd177, %rd163, %rd149; | |
| mul.f32x2 %rd178, %rd163, %rd148; | |
| mul.f32x2 %rd179, %rd163, %rd147; | |
| .loc 1 189 31 // bar.py:189:31 | |
| fma.rn.f32x2 %rd233, %rd146, %rd179, %rd233; | |
| fma.rn.f32x2 %rd232, %rd145, %rd178, %rd232; | |
| fma.rn.f32x2 %rd231, %rd144, %rd177, %rd231; | |
| fma.rn.f32x2 %rd230, %rd143, %rd176, %rd230; | |
| fma.rn.f32x2 %rd229, %rd142, %rd175, %rd229; | |
| fma.rn.f32x2 %rd228, %rd141, %rd174, %rd228; | |
| fma.rn.f32x2 %rd227, %rd140, %rd173, %rd227; | |
| fma.rn.f32x2 %rd226, %rd139, %rd172, %rd226; | |
| fma.rn.f32x2 %rd225, %rd138, %rd171, %rd225; | |
| fma.rn.f32x2 %rd224, %rd137, %rd170, %rd224; | |
| fma.rn.f32x2 %rd223, %rd136, %rd169, %rd223; | |
| fma.rn.f32x2 %rd222, %rd135, %rd168, %rd222; | |
| fma.rn.f32x2 %rd221, %rd134, %rd167, %rd221; | |
| fma.rn.f32x2 %rd220, %rd133, %rd166, %rd220; | |
| fma.rn.f32x2 %rd234, %rd131, %rd165, %rd234; | |
| fma.rn.f32x2 %rd235, %rd132, %rd164, %rd235; | |
| $L__BB0_21: | |
| .loc 1 0 31 // bar.py:0:31 | |
| cvt.u32.u64 %r265, %rd6; | |
| cvt.u32.u64 %r266, %rd5; | |
| cvt.u32.u64 %r267, %rd4; | |
| cvt.u32.u64 %r268, %rd3; | |
| .loc 1 114 30 // bar.py:114:30 | |
| setp.lt.s32 %p40, %r265, %r21; | |
| setp.lt.s32 %p41, %r266, %r21; | |
| setp.lt.s32 %p42, %r267, %r21; | |
| setp.lt.s32 %p43, %r268, %r21; | |
| .loc 1 172 22 // bar.py:172:22 | |
| cp.async.wait_group 0; | |
| bar.sync 0; | |
| elect.sync %r269|%p44, -1; | |
| and.pred %p34, %p9, %p44; | |
| add.s32 %r231, %r29, 12544; | |
| // begin inline asm | |
| @%p34 mbarrier.inval.shared::cta.b64 [%r231]; | |
| // end inline asm | |
| elect.sync %r270|%p45, -1; | |
| and.pred %p35, %p9, %p45; | |
| // begin inline asm | |
| @%p35 mbarrier.inval.shared::cta.b64 [%r71]; | |
| // end inline asm | |
| .loc 1 212 33 // bar.py:212:33 | |
| mov.b64 {%r271, %r272}, %rd220; | |
| cvt.rn.bf16x2.f32 %r233, %r272, %r271; | |
| mov.b64 {%r273, %r274}, %rd221; | |
| cvt.rn.bf16x2.f32 %r237, %r274, %r273; | |
| mov.b64 {%r275, %r276}, %rd222; | |
| cvt.rn.bf16x2.f32 %r241, %r276, %r275; | |
| mov.b64 {%r277, %r278}, %rd223; | |
| cvt.rn.bf16x2.f32 %r245, %r278, %r277; | |
| mov.b64 {%r279, %r280}, %rd224; | |
| cvt.rn.bf16x2.f32 %r234, %r280, %r279; | |
| mov.b64 {%r281, %r282}, %rd225; | |
| cvt.rn.bf16x2.f32 %r238, %r282, %r281; | |
| mov.b64 {%r283, %r284}, %rd226; | |
| cvt.rn.bf16x2.f32 %r242, %r284, %r283; | |
| mov.b64 {%r285, %r286}, %rd227; | |
| cvt.rn.bf16x2.f32 %r246, %r286, %r285; | |
| mov.b64 {%r287, %r288}, %rd228; | |
| cvt.rn.bf16x2.f32 %r235, %r288, %r287; | |
| mov.b64 {%r289, %r290}, %rd229; | |
| cvt.rn.bf16x2.f32 %r239, %r290, %r289; | |
| mov.b64 {%r291, %r292}, %rd230; | |
| cvt.rn.bf16x2.f32 %r243, %r292, %r291; | |
| mov.b64 {%r293, %r294}, %rd231; | |
| cvt.rn.bf16x2.f32 %r247, %r294, %r293; | |
| mov.b64 {%r295, %r296}, %rd232; | |
| cvt.rn.bf16x2.f32 %r236, %r296, %r295; | |
| mov.b64 {%r297, %r298}, %rd233; | |
| cvt.rn.bf16x2.f32 %r240, %r298, %r297; | |
| mov.b64 {%r299, %r300}, %rd234; | |
| cvt.rn.bf16x2.f32 %r244, %r300, %r299; | |
| mov.b64 {%r301, %r302}, %rd235; | |
| cvt.rn.bf16x2.f32 %r248, %r302, %r301; | |
| .loc 1 213 37 // bar.py:213:37 | |
| or.b32 %r303, %r10, %r8; | |
| .loc 1 214 33 // bar.py:214:33 | |
| cvt.s64.s32 %rd184, %r25; | |
| mul.lo.s64 %rd185, %rd3, %rd184; | |
| mul.lo.s64 %rd186, %rd4, %rd184; | |
| mul.lo.s64 %rd187, %rd5, %rd184; | |
| mul.lo.s64 %rd188, %rd6, %rd184; | |
| .loc 1 214 21 // bar.py:214:21 | |
| shl.b64 %rd189, %rd185, 1; | |
| add.s64 %rd190, %rd22, %rd189; | |
| shl.b64 %rd191, %rd186, 1; | |
| add.s64 %rd192, %rd22, %rd191; | |
| shl.b64 %rd193, %rd187, 1; | |
| add.s64 %rd194, %rd22, %rd193; | |
| shl.b64 %rd195, %rd188, 1; | |
| add.s64 %rd196, %rd22, %rd195; | |
| .loc 1 214 55 // bar.py:214:55 | |
| mul.wide.s32 %rd197, %r303, 2; | |
| add.s64 %rd180, %rd190, %rd197; | |
| add.s64 %rd181, %rd192, %rd197; | |
| add.s64 %rd182, %rd194, %rd197; | |
| add.s64 %rd183, %rd196, %rd197; | |
| .loc 1 215 55 // bar.py:215:55 | |
| setp.lt.s32 %p46, %r303, %r19; | |
| .loc 1 215 36 // bar.py:215:36 | |
| and.pred %p36, %p46, %p43; | |
| and.pred %p37, %p46, %p42; | |
| and.pred %p38, %p46, %p41; | |
| and.pred %p39, %p46, %p40; | |
| .loc 1 216 21 // bar.py:216:21 | |
| and.b32 %r304, %r14, 176; | |
| shl.b32 %r305, %r6, 3; | |
| bfe.s32 %r306, %r1, 2, 1; | |
| and.b32 %r307, %r306, 1088; | |
| shl.b32 %r308, %r1, 2; | |
| and.b32 %r309, %r308, 64; | |
| xor.b32 %r310, %r307, %r309; | |
| add.s32 %r311, %r29, %r304; | |
| add.s32 %r312, %r311, %r305; | |
| add.s32 %r232, %r312, %r310; | |
| // begin inline asm | |
| st.shared.v4.b32 [ %r232 + 0 ], { %r233, %r234, %r235, %r236 }; | |
| // end inline asm | |
| bar.sync 0; | |
| shl.b32 %r313, %r1, 5; | |
| and.b32 %r314, %r313, 1792; | |
| and.b32 %r315, %r7, 48; | |
| shl.b32 %r316, %r6, 1; | |
| shl.b32 %r317, %r11, 6; | |
| xor.b32 %r318, %r316, %r317; | |
| add.s32 %r319, %r29, %r314; | |
| add.s32 %r320, %r319, %r315; | |
| add.s32 %r321, %r320, %r318; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r249, %r253, %r257, %r261}, [%r321]; | |
| bar.sync 0; | |
| // begin inline asm | |
| st.shared.v4.b32 [ %r232 + 0 ], { %r237, %r238, %r239, %r240 }; | |
| // end inline asm | |
| bar.sync 0; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r250, %r254, %r258, %r262}, [%r321]; | |
| bar.sync 0; | |
| // begin inline asm | |
| st.shared.v4.b32 [ %r232 + 0 ], { %r241, %r242, %r243, %r244 }; | |
| // end inline asm | |
| bar.sync 0; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r251, %r255, %r259, %r263}, [%r321]; | |
| bar.sync 0; | |
| // begin inline asm | |
| st.shared.v4.b32 [ %r232 + 0 ], { %r245, %r246, %r247, %r248 }; | |
| // end inline asm | |
| bar.sync 0; | |
| ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r252, %r256, %r260, %r264}, [%r321]; | |
| // begin inline asm | |
| @%p36 st.global.v4.b32 [ %rd180 + 0 ], { %r249, %r250, %r251, %r252 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p37 st.global.v4.b32 [ %rd181 + 0 ], { %r253, %r254, %r255, %r256 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p38 st.global.v4.b32 [ %rd182 + 0 ], { %r257, %r258, %r259, %r260 }; | |
| // end inline asm | |
| // begin inline asm | |
| @%p39 st.global.v4.b32 [ %rd183 + 0 ], { %r261, %r262, %r263, %r264 }; | |
| // end inline asm | |
| .loc 1 216 4 // bar.py:216:4 | |
| bar.sync 0; | |
| // begin inline asm | |
| @%p1 tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r2, 64; | |
| // end inline asm | |
| $L__BB0_2: // %common.ret | |
| .loc 1 0 0 // bar.py:0 | |
| ret; | |
| $L__tmp8: | |
| $L__func_end0: | |
| // -- End function | |
| } | |
| .file 1 "/home/dev/bar.py" | |
| .file 2 "/home/dev/py3.12-torch2.12/lib/python3.12/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 0 // EOM(3) | |
| } | |
| .section .debug_info | |
| { | |
| .b32 176 // 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:0xa9 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 98 // DW_AT_name | |
| .b8 97 | |
| .b8 114 | |
| .b8 46 | |
| .b8 112 | |
| .b8 121 | |
| .b8 0 | |
| .b32 .debug_line // DW_AT_stmt_list | |
| .b8 47 // DW_AT_comp_dir | |
| .b8 104 | |
| .b8 111 | |
| .b8 109 | |
| .b8 101 | |
| .b8 47 | |
| .b8 100 | |
| .b8 101 | |
| .b8 118 | |
| .b8 0 | |
| .b8 2 // Abbrev [2] 0x2a:0x13 DW_TAG_subprogram | |
| .b8 102 // DW_AT_name | |
| .b8 117 | |
| .b8 115 | |
| .b8 101 | |
| .b8 100 | |
| .b8 95 | |
| .b8 109 | |
| .b8 111 | |
| .b8 101 | |
| .b8 95 | |
| .b8 107 | |
| .b8 101 | |
| .b8 114 | |
| .b8 110 | |
| .b8 101 | |
| .b8 108 | |
| .b8 0 | |
| .b8 1 // DW_AT_inline | |
| .b8 3 // Abbrev [3] 0x3d:0x76 DW_TAG_subprogram | |
| .b64 $L__func_begin0 // DW_AT_low_pc | |
| .b64 $L__func_end0 // DW_AT_high_pc | |
| .b32 42 // DW_AT_abstract_origin | |
| .b8 4 // Abbrev [4] 0x52:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // 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 94 // DW_AT_call_line | |
| .b8 28 // DW_AT_call_column | |
| .b8 4 // Abbrev [4] 0x6a:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // DW_AT_abstract_origin | |
| .b64 $L__tmp2 // DW_AT_low_pc | |
| .b64 $L__tmp3 // DW_AT_high_pc | |
| .b8 1 // DW_AT_call_file | |
| .b8 95 // DW_AT_call_line | |
| .b8 27 // DW_AT_call_column | |
| .b8 4 // Abbrev [4] 0x82:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // DW_AT_abstract_origin | |
| .b64 $L__tmp4 // DW_AT_low_pc | |
| .b64 $L__tmp5 // DW_AT_high_pc | |
| .b8 1 // DW_AT_call_file | |
| .b8 128 // DW_AT_call_line | |
| .b8 12 // DW_AT_call_column | |
| .b8 4 // Abbrev [4] 0x9a:0x18 DW_TAG_inlined_subroutine | |
| .b32 42 // DW_AT_abstract_origin | |
| .b64 $L__tmp6 // DW_AT_low_pc | |
| .b64 $L__tmp7 // DW_AT_high_pc | |
| .b8 1 // DW_AT_call_file | |
| .b8 172 // DW_AT_call_line | |
| .b8 33 // DW_AT_call_column | |
| .b8 0 // End Of Children Mark | |
| .b8 0 // End Of Children Mark | |
| } | |
| .section .debug_macinfo { } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| --- kernel_37.ptx 2026-06-05 15:04:59 | |
| +++ kernel_37_fixed.ptx 2026-06-05 15:04:58 | |
| @@ -703,6 +703,7 @@ | |
| mov.b64 %rd103, {%r131, %r132}; | |
| mov.b64 %rd104, {%r161, %r162}; | |
| tcgen05.wait::ld.sync.aligned; | |
| + fence.proxy.async; | |
| .loc 1 189 46 // bar.py:189:46 | |
| mov.b64 %rd105, {%r129, %r129}; | |
| mul.f32x2 %rd106, %rd105, %rd104; |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| """ | |
| Compile PTX to cubin and test determinism. | |
| Usage: python test_ptx.py <ptx> [ptx2] | |
| """ | |
| import ctypes, subprocess, sys, tempfile | |
| import torch, triton | |
| from triton.runtime import driver | |
| E, M, K, N, TOPK = 128, 1958, 2048, 1536, 8 | |
| BLOCK_M, BLOCK_N, SMEM = 64, 64, 12560 | |
| libcuda = ctypes.CDLL("libcuda.so.1") | |
| def make_inputs(): | |
| r = torch.arange(M, device="cuda").reshape(M, 1) | |
| c = torch.arange(K, device="cuda").reshape(1, K) | |
| A = (((r + c) % 17) - 8).to(torch.bfloat16).div(80).to(torch.float8_e4m3fn).contiguous() | |
| e = torch.arange(E, device="cuda").reshape(E, 1, 1) | |
| n = torch.arange(N, device="cuda").reshape(1, N, 1) | |
| k = torch.arange(K, device="cuda").reshape(1, 1, K) | |
| B = (((e + n + k) % 19) - 9).to(torch.bfloat16).div(90).to(torch.float8_e4m3fn).contiguous() | |
| A_scale = torch.full((M, 16), 0.01, dtype=torch.float32, device="cuda") | |
| B_scale = torch.full((E, 12, 16), 0.01, dtype=torch.float32, device="cuda") | |
| topk_ids = torch.tensor([69, 10, 64, 77, 85, 15, 75, 87], | |
| device="cuda", dtype=torch.int32).repeat(M, 1).contiguous() | |
| # align_block_size | |
| flat = topk_ids.reshape(-1) | |
| num_tokens = flat.numel() | |
| sorted_chunks, expert_chunks = [], [] | |
| for exp in range(E): | |
| tids = torch.nonzero(flat == exp, as_tuple=False).flatten().flip(0) | |
| if tids.numel() == 0: | |
| continue | |
| pad = (-tids.numel()) % BLOCK_M | |
| if pad: | |
| tids = torch.cat([tids, torch.full((pad,), num_tokens, device="cuda", dtype=tids.dtype)]) | |
| sorted_chunks.append(tids.to(torch.int32)) | |
| expert_chunks.append(torch.full((tids.numel() // BLOCK_M,), exp, device="cuda", dtype=torch.int32)) | |
| sorted_ids = torch.cat(sorted_chunks).contiguous() | |
| expert_ids = torch.cat(expert_chunks).contiguous() | |
| ntpp = torch.tensor([sorted_ids.numel()], device="cuda", dtype=torch.int32) | |
| return A, B, A_scale, B_scale, sorted_ids, expert_ids, ntpp, num_tokens | |
| def launch(func, A, B, C, A_scale, B_scale, sorted_ids, expert_ids, ntpp, EM, num_tokens): | |
| d = driver.active | |
| stream = d.get_current_stream(d.get_current_device()) | |
| grid_x = triton.cdiv(EM, BLOCK_M) * triton.cdiv(N, BLOCK_N) | |
| params = [ | |
| ("Q", A.data_ptr()), ("Q", B.data_ptr()), ("Q", C.data_ptr()), | |
| ("Q", A_scale.data_ptr()), ("Q", B_scale.data_ptr()), | |
| ("Q", sorted_ids.data_ptr()), ("Q", expert_ids.data_ptr()), ("Q", ntpp.data_ptr()), | |
| ("I", N), ("I", K), ("I", EM), ("I", num_tokens), | |
| ("I", A.stride(0)), ("I", B.stride(0)), ("I", B.stride(1)), ("I", C.stride(1)), | |
| ("I", A_scale.stride(0)), ("I", B_scale.stride(0)), ("I", B_scale.stride(1)), | |
| ("I", 0), ("I", 0), ("Q", 0), ("Q", 0), | |
| ] | |
| bufs = [] | |
| ptrs = (ctypes.c_void_p * len(params))() | |
| for i, (fmt, val) in enumerate(params): | |
| b = ctypes.c_uint64(val) if fmt == "Q" else ctypes.c_uint32(val) | |
| bufs.append(b) | |
| ptrs[i] = ctypes.cast(ctypes.pointer(b), ctypes.c_void_p) | |
| libcuda.cuLaunchKernel(ctypes.c_void_p(func), grid_x, 1, 1, 128, 1, 1, | |
| SMEM, ctypes.c_void_p(stream), ptrs, None) | |
| torch.cuda.synchronize() | |
| def test(ptx_path, trials=10): | |
| print(f"Testing: {ptx_path}") | |
| with tempfile.NamedTemporaryFile(suffix=".cubin") as f: | |
| subprocess.check_call(["ptxas", "--gpu-name", "sm_100a", "-o", f.name, ptx_path], | |
| stderr=subprocess.PIPE) | |
| cubin = f.read() | |
| d = driver.active | |
| _, func, *_ = d.utils.load_binary("fused_moe_kernel", cubin, SMEM, d.get_current_device()) | |
| A, B, A_scale, B_scale, sorted_ids, expert_ids, ntpp, num_tokens = make_inputs() | |
| EM = sorted_ids.size(0) | |
| results = [] | |
| for _ in range(trials): | |
| C = torch.empty((M, TOPK, N), dtype=torch.bfloat16, device="cuda") | |
| launch(func, A, B, C, A_scale, B_scale, sorted_ids, expert_ids, ntpp, EM, num_tokens) | |
| results.append(C.clone()) | |
| diffs = sum(1 for i in range(trials - 1) if not torch.equal(results[i], results[i + 1])) | |
| if diffs == 0: | |
| print(f" DETERMINISTIC ({trials} runs)") | |
| else: | |
| md = max((results[i].float() - results[i+1].float()).abs().max().item() | |
| for i in range(trials - 1) if not torch.equal(results[i], results[i+1])) | |
| print(f" NON-DETERMINISTIC ({diffs}/{trials-1} pairs differ, max_diff={md:.2e})") | |
| return results[0], diffs == 0 | |
| if __name__ == "__main__": | |
| torch.set_default_device("cuda") | |
| if len(sys.argv) < 2: | |
| print("Usage: python test_ptx.py <ptx> [ptx2]") | |
| sys.exit(1) | |
| r1, d1 = test(sys.argv[1]) | |
| if len(sys.argv) > 2: | |
| r2, d2 = test(sys.argv[2]) | |
| if d1 and d2: | |
| if torch.equal(r1, r2): | |
| print("\nBoth deterministic and MATCH") | |
| else: | |
| diff = (r1.float() - r2.float()).abs() | |
| print(f"\nBoth deterministic but DIFFER: max={diff.max().item():.2e}") |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment