| // | |
| // Generated by LLVM NVPTX Back-End | |
| // | |
| .version 8.2 | |
| .target sm_89 | |
| .address_size 64 | |
| // .globl triton__0d1d2d3d4d5de6de | |
| .extern .func __assertfail | |
| ( | |
| .param .b64 __assertfail_param_0, | |
| .param .b64 __assertfail_param_1, | |
| .param .b32 __assertfail_param_2, | |
| .param .b64 __assertfail_param_3, | |
| .param .b64 __assertfail_param_4 | |
| ) | |
| ; | |
| .global .align 1 .b8 assertFunc_1[25] = {95, 99, 97, 108, 108, 95, 119, 105, 116, 104, 95, 102, 114, 97, 109, 101, 115, 95, 114, 101, 109, 111, 118, 101, 100}; | |
| .global .align 1 .b8 assertFile_1[38] = {60, 102, 114, 111, 122, 101, 110, 32, 105, 109, 112, 111, 114, 116, 108, 105, 98, 46, 95, 98, 111, 111, 116, 115, 116, 114, 97, 112, 95, 101, 120, 116, 101, 114, 110, 97, 108, 62}; | |
| .global .align 1 .b8 assertMessage_1[39] = {105, 110, 100, 101, 120, 32, 111, 117, 116, 32, 111, 102, 32, 98, 111, 117, 110, 100, 115, 58, 32, 48, 32, 60, 61, 32, 116, 109, 112, 49, 51, 32, 60, 32, 53, 48, 50, 53, 55}; | |
| .global .align 1 .b8 assertFunc_0[25] = {95, 99, 97, 108, 108, 95, 119, 105, 116, 104, 95, 102, 114, 97, 109, 101, 115, 95, 114, 101, 109, 111, 118, 101, 100}; | |
| .global .align 1 .b8 assertFile_0[38] = {60, 102, 114, 111, 122, 101, 110, 32, 105, 109, 112, 111, 114, 116, 108, 105, 98, 46, 95, 98, 111, 111, 116, 115, 116, 114, 97, 112, 95, 101, 120, 116, 101, 114, 110, 97, 108, 62}; | |
| .global .align 1 .b8 assertMessage_0[38] = {105, 110, 100, 101, 120, 32, 111, 117, 116, 32, 111, 102, 32, 98, 111, 117, 110, 100, 115, 58, 32, 48, 32, 60, 61, 32, 116, 109, 112, 51, 32, 60, 32, 53, 48, 50, 53, 55}; | |
| .extern .shared .align 1 .b8 global_smem[]; | |
| .global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0}; | |
| .visible .entry triton__0d1d2d3d4d5de6de( | |
| .param .u64 triton__0d1d2d3d4d5de6de_param_0, | |
| .param .u64 triton__0d1d2d3d4d5de6de_param_1, | |
| .param .u64 triton__0d1d2d3d4d5de6de_param_2, | |
| .param .u64 triton__0d1d2d3d4d5de6de_param_3, | |
| .param .u64 triton__0d1d2d3d4d5de6de_param_4, | |
| .param .u32 triton__0d1d2d3d4d5de6de_param_5, | |
| .param .u32 triton__0d1d2d3d4d5de6de_param_6 | |
| ) | |
| .maxntid 256, 1, 1 | |
| { | |
| .reg .pred %p<137>; | |
| .reg .b16 %rs<17>; | |
| .reg .b32 %r<408>; | |
| .reg .f32 %f<614>; | |
| .reg .b64 %rd<107>; | |
| .loc 1 18 0 | |
| $L__func_begin0: | |
| .loc 1 18 0 | |
| ld.param.u64 %rd13, [triton__0d1d2d3d4d5de6de_param_4]; | |
| ld.param.u64 %rd12, [triton__0d1d2d3d4d5de6de_param_3]; | |
| ld.param.u64 %rd49, [triton__0d1d2d3d4d5de6de_param_0]; | |
| ld.param.u64 %rd50, [triton__0d1d2d3d4d5de6de_param_1]; | |
| $L__tmp0: | |
| .loc 1 22 44 | |
| mov.u32 %r13, %tid.x; | |
| ld.param.u64 %rd51, [triton__0d1d2d3d4d5de6de_param_2]; | |
| bfe.u32 %r1, %r13, 3, 5; | |
| and.b32 %r2, %r13, 63; | |
| .loc 1 24 33 | |
| shl.b32 %r14, %r13, 3; | |
| and.b32 %r3, %r14, 56; | |
| .loc 1 31 36 | |
| shr.u32 %r4, %r13, 6; | |
| .loc 1 21 28 | |
| mov.u32 %r11, %ctaid.x; | |
| .loc 1 21 33 | |
| shl.b32 %r15, %r11, 6; | |
| .loc 1 22 23 | |
| or.b32 %r16, %r15, %r1; | |
| or.b32 %r17, %r16, 32; | |
| or.b32 %r18, %r15, %r2; | |
| .loc 1 26 30 | |
| mul.wide.s32 %rd52, %r16, 8; | |
| add.s64 %rd15, %rd49, %rd52; | |
| add.s64 %rd31, %rd15, 256; | |
| mul.wide.s32 %rd53, %r18, 8; | |
| add.s64 %rd47, %rd49, %rd53; | |
| mov.pred %p1, -1; | |
| .loc 1 26 35 | |
| mov.u64 %rd14, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd14 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd16, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd16 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd18, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd18 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd20, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd20 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd22, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd22 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd24, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd24 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd26, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd26 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd28, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd28 }, [ %rd15 + 0 ]; | |
| mov.u64 %rd30, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd30 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd32, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd32 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd34, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd34 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd36, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd36 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd38, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd38 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd40, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd40 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd42, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd42 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd44, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd44 }, [ %rd31 + 0 ]; | |
| mov.u64 %rd46, 0x0; | |
| @%p1 ld.global.L1::evict_last.b64 { %rd46 }, [ %rd47 + 0 ]; | |
| .loc 1 27 18 | |
| bfe.s32 %r19, %r11, 25, 1; | |
| shr.u32 %r20, %r19, 23; | |
| add.s32 %r21, %r16, %r20; | |
| and.b32 %r22, %r21, 16776704; | |
| sub.s32 %r23, %r16, %r22; | |
| add.s32 %r24, %r17, %r20; | |
| and.b32 %r25, %r24, 16776704; | |
| sub.s32 %r26, %r17, %r25; | |
| .loc 1 35 44 | |
| shl.b32 %r27, %r23, 8; | |
| shl.b32 %r28, %r26, 8; | |
| .loc 1 36 22 | |
| add.s64 %rd54, %rd46, 50257; | |
| .loc 1 37 22 | |
| setp.lt.s64 %p18, %rd14, 0; | |
| setp.lt.s64 %p19, %rd30, 0; | |
| setp.lt.s64 %p20, %rd46, 0; | |
| .loc 1 38 36 | |
| selp.b64 %rd1, %rd54, %rd46, %p20; | |
| .loc 1 40 44 | |
| shl.b64 %rd55, %rd14, 8; | |
| add.s64 %rd56, %rd55, 12865792; | |
| selp.b64 %rd57, %rd56, %rd55, %p18; | |
| shl.b64 %rd58, %rd30, 8; | |
| add.s64 %rd59, %rd58, 12865792; | |
| selp.b64 %rd60, %rd59, %rd58, %p19; | |
| .loc 1 31 36 | |
| and.b32 %r29, %r13, 7; | |
| mul.wide.u32 %rd2, %r29, 32; | |
| shl.b64 %rd61, %rd60, 2; | |
| or.b64 %rd62, %rd2, %rd61; | |
| add.s64 %rd3, %rd50, %rd62; | |
| shl.b64 %rd63, %rd57, 2; | |
| or.b64 %rd64, %rd2, %rd63; | |
| add.s64 %rd4, %rd50, %rd64; | |
| or.b32 %r30, %r28, %r3; | |
| mul.wide.s32 %rd65, %r30, 4; | |
| add.s64 %rd5, %rd51, %rd65; | |
| or.b32 %r31, %r27, %r3; | |
| mul.wide.s32 %rd66, %r31, 4; | |
| add.s64 %rd6, %rd51, %rd66; | |
| mov.f32 %f550, 0f00000000; | |
| mov.u64 %rd105, 0; | |
| mov.b32 %r406, -64; | |
| mov.f32 %f551, %f550; | |
| mov.f32 %f552, %f550; | |
| mov.f32 %f553, %f550; | |
| mov.f32 %f554, %f550; | |
| mov.f32 %f555, %f550; | |
| mov.f32 %f556, %f550; | |
| mov.f32 %f557, %f550; | |
| mov.f32 %f558, %f550; | |
| mov.f32 %f559, %f550; | |
| mov.f32 %f560, %f550; | |
| mov.f32 %f561, %f550; | |
| mov.f32 %f562, %f550; | |
| mov.f32 %f563, %f550; | |
| mov.f32 %f564, %f550; | |
| mov.f32 %f565, %f550; | |
| mov.f32 %f566, %f550; | |
| mov.f32 %f567, %f550; | |
| mov.f32 %f568, %f550; | |
| mov.f32 %f569, %f550; | |
| mov.f32 %f570, %f550; | |
| mov.f32 %f571, %f550; | |
| mov.f32 %f572, %f550; | |
| mov.f32 %f573, %f550; | |
| mov.f32 %f574, %f550; | |
| mov.f32 %f575, %f550; | |
| mov.f32 %f576, %f550; | |
| mov.f32 %f577, %f550; | |
| mov.f32 %f578, %f550; | |
| mov.f32 %f579, %f550; | |
| mov.f32 %f580, %f550; | |
| mov.f32 %f581, %f550; | |
| mov.f32 %f582, %f550; | |
| mov.f32 %f583, %f550; | |
| mov.f32 %f584, %f550; | |
| mov.f32 %f585, %f550; | |
| mov.f32 %f586, %f550; | |
| mov.f32 %f587, %f550; | |
| mov.f32 %f588, %f550; | |
| mov.f32 %f589, %f550; | |
| mov.f32 %f590, %f550; | |
| mov.f32 %f591, %f550; | |
| mov.f32 %f592, %f550; | |
| mov.f32 %f593, %f550; | |
| mov.f32 %f594, %f550; | |
| mov.f32 %f595, %f550; | |
| mov.f32 %f596, %f550; | |
| mov.f32 %f597, %f550; | |
| mov.f32 %f598, %f550; | |
| mov.f32 %f599, %f550; | |
| mov.f32 %f600, %f550; | |
| mov.f32 %f601, %f550; | |
| mov.f32 %f602, %f550; | |
| mov.f32 %f603, %f550; | |
| mov.f32 %f604, %f550; | |
| mov.f32 %f605, %f550; | |
| mov.f32 %f606, %f550; | |
| mov.f32 %f607, %f550; | |
| mov.f32 %f608, %f550; | |
| mov.f32 %f609, %f550; | |
| mov.f32 %f610, %f550; | |
| mov.f32 %f611, %f550; | |
| mov.f32 %f612, %f550; | |
| mov.f32 %f613, %f550; | |
| bra.uni $L__BB0_1; | |
| $L__BB0_3: | |
| .loc 1 40 40 | |
| add.s64 %rd78, %rd4, %rd105; | |
| .loc 1 40 34 | |
| add.s64 %rd79, %rd78, 16; | |
| add.s64 %rd80, %rd3, %rd105; | |
| .loc 1 40 52 | |
| add.s64 %rd81, %rd80, 16; | |
| mov.u32 %r65, 0x0; | |
| mov.u32 %r66, 0x0; | |
| mov.u32 %r67, 0x0; | |
| mov.u32 %r68, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r65, %r66, %r67, %r68 }, [ %rd78 + 0 ]; | |
| @!%p1 mov.u32 %r65, %r342; | |
| @!%p1 mov.u32 %r66, %r342; | |
| @!%p1 mov.u32 %r67, %r342; | |
| @!%p1 mov.u32 %r68, %r342; | |
| mov.b32 %f174, %r65; | |
| mov.b32 %f175, %r66; | |
| mov.b32 %f176, %r67; | |
| mov.b32 %f177, %r68; | |
| mov.u32 %r73, 0x0; | |
| mov.u32 %r74, 0x0; | |
| mov.u32 %r75, 0x0; | |
| mov.u32 %r76, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r73, %r74, %r75, %r76 }, [ %rd79 + 0 ]; | |
| @!%p1 mov.u32 %r73, %r342; | |
| @!%p1 mov.u32 %r74, %r342; | |
| @!%p1 mov.u32 %r75, %r342; | |
| @!%p1 mov.u32 %r76, %r342; | |
| mov.b32 %f178, %r73; | |
| mov.b32 %f179, %r74; | |
| mov.b32 %f180, %r75; | |
| mov.b32 %f181, %r76; | |
| mov.u32 %r81, 0x0; | |
| mov.u32 %r82, 0x0; | |
| mov.u32 %r83, 0x0; | |
| mov.u32 %r84, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r81, %r82, %r83, %r84 }, [ %rd80 + 0 ]; | |
| @!%p1 mov.u32 %r81, %r342; | |
| @!%p1 mov.u32 %r82, %r342; | |
| @!%p1 mov.u32 %r83, %r342; | |
| @!%p1 mov.u32 %r84, %r342; | |
| mov.b32 %f182, %r81; | |
| mov.b32 %f183, %r82; | |
| mov.b32 %f184, %r83; | |
| mov.b32 %f185, %r84; | |
| mov.u32 %r89, 0x0; | |
| mov.u32 %r90, 0x0; | |
| mov.u32 %r91, 0x0; | |
| mov.u32 %r92, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r89, %r90, %r91, %r92 }, [ %rd81 + 0 ]; | |
| @!%p1 mov.u32 %r89, %r342; | |
| @!%p1 mov.u32 %r90, %r342; | |
| @!%p1 mov.u32 %r91, %r342; | |
| @!%p1 mov.u32 %r92, %r342; | |
| mov.b32 %f186, %r89; | |
| mov.b32 %f187, %r90; | |
| mov.b32 %f188, %r91; | |
| mov.b32 %f189, %r92; | |
| .loc 1 41 22 | |
| add.f32 %f190, %f65, %f174; | |
| add.f32 %f191, %f66, %f175; | |
| add.f32 %f192, %f67, %f176; | |
| add.f32 %f193, %f68, %f177; | |
| add.f32 %f194, %f69, %f178; | |
| add.f32 %f195, %f70, %f179; | |
| add.f32 %f196, %f71, %f180; | |
| add.f32 %f197, %f72, %f181; | |
| add.f32 %f198, %f73, %f182; | |
| add.f32 %f199, %f74, %f183; | |
| add.f32 %f200, %f75, %f184; | |
| add.f32 %f201, %f76, %f185; | |
| add.f32 %f202, %f77, %f186; | |
| add.f32 %f203, %f78, %f187; | |
| add.f32 %f204, %f79, %f188; | |
| add.f32 %f205, %f80, %f189; | |
| $L__tmp1: | |
| .loc 2 96 20 | |
| sub.f32 %f206, %f190, %f598; | |
| sub.f32 %f207, %f191, %f599; | |
| sub.f32 %f208, %f192, %f600; | |
| sub.f32 %f209, %f193, %f601; | |
| sub.f32 %f210, %f194, %f602; | |
| sub.f32 %f211, %f195, %f603; | |
| sub.f32 %f212, %f196, %f604; | |
| sub.f32 %f213, %f197, %f605; | |
| sub.f32 %f214, %f198, %f606; | |
| sub.f32 %f215, %f199, %f607; | |
| sub.f32 %f216, %f200, %f608; | |
| sub.f32 %f217, %f201, %f609; | |
| sub.f32 %f218, %f202, %f610; | |
| sub.f32 %f219, %f203, %f611; | |
| sub.f32 %f220, %f204, %f612; | |
| sub.f32 %f221, %f205, %f613; | |
| .loc 2 97 26 | |
| add.f32 %f550, %f550, 0f3F800000; | |
| add.f32 %f551, %f551, 0f3F800000; | |
| add.f32 %f552, %f552, 0f3F800000; | |
| add.f32 %f553, %f553, 0f3F800000; | |
| add.f32 %f554, %f554, 0f3F800000; | |
| add.f32 %f555, %f555, 0f3F800000; | |
| add.f32 %f556, %f556, 0f3F800000; | |
| add.f32 %f557, %f557, 0f3F800000; | |
| add.f32 %f558, %f558, 0f3F800000; | |
| add.f32 %f559, %f559, 0f3F800000; | |
| add.f32 %f560, %f560, 0f3F800000; | |
| add.f32 %f561, %f561, 0f3F800000; | |
| add.f32 %f562, %f562, 0f3F800000; | |
| add.f32 %f563, %f563, 0f3F800000; | |
| add.f32 %f564, %f564, 0f3F800000; | |
| add.f32 %f565, %f565, 0f3F800000; | |
| add.f32 %f566, %f566, 0f3F800000; | |
| add.f32 %f567, %f567, 0f3F800000; | |
| add.f32 %f568, %f568, 0f3F800000; | |
| add.f32 %f569, %f569, 0f3F800000; | |
| add.f32 %f570, %f570, 0f3F800000; | |
| add.f32 %f571, %f571, 0f3F800000; | |
| add.f32 %f572, %f572, 0f3F800000; | |
| add.f32 %f573, %f573, 0f3F800000; | |
| add.f32 %f574, %f574, 0f3F800000; | |
| add.f32 %f575, %f575, 0f3F800000; | |
| add.f32 %f576, %f576, 0f3F800000; | |
| add.f32 %f577, %f577, 0f3F800000; | |
| add.f32 %f578, %f578, 0f3F800000; | |
| add.f32 %f579, %f579, 0f3F800000; | |
| add.f32 %f580, %f580, 0f3F800000; | |
| add.f32 %f581, %f581, 0f3F800000; | |
| .loc 2 98 30 | |
| mov.b32 %r98, %f206; | |
| mov.b32 %r99, %f550; | |
| div.full.f32 %r97, %r98, %r99; | |
| mov.b32 %f222, %r97; | |
| mov.b32 %r101, %f207; | |
| mov.b32 %r102, %f551; | |
| div.full.f32 %r100, %r101, %r102; | |
| mov.b32 %f223, %r100; | |
| mov.b32 %r104, %f208; | |
| mov.b32 %r105, %f552; | |
| div.full.f32 %r103, %r104, %r105; | |
| mov.b32 %f224, %r103; | |
| mov.b32 %r107, %f209; | |
| mov.b32 %r108, %f553; | |
| div.full.f32 %r106, %r107, %r108; | |
| mov.b32 %f225, %r106; | |
| mov.b32 %r110, %f210; | |
| mov.b32 %r111, %f554; | |
| div.full.f32 %r109, %r110, %r111; | |
| mov.b32 %f226, %r109; | |
| mov.b32 %r113, %f211; | |
| mov.b32 %r114, %f555; | |
| div.full.f32 %r112, %r113, %r114; | |
| mov.b32 %f227, %r112; | |
| mov.b32 %r116, %f212; | |
| mov.b32 %r117, %f556; | |
| div.full.f32 %r115, %r116, %r117; | |
| mov.b32 %f228, %r115; | |
| mov.b32 %r119, %f213; | |
| mov.b32 %r120, %f557; | |
| div.full.f32 %r118, %r119, %r120; | |
| mov.b32 %f229, %r118; | |
| mov.b32 %r122, %f214; | |
| mov.b32 %r123, %f558; | |
| div.full.f32 %r121, %r122, %r123; | |
| mov.b32 %f230, %r121; | |
| mov.b32 %r125, %f215; | |
| mov.b32 %r126, %f559; | |
| div.full.f32 %r124, %r125, %r126; | |
| mov.b32 %f231, %r124; | |
| mov.b32 %r128, %f216; | |
| mov.b32 %r129, %f560; | |
| div.full.f32 %r127, %r128, %r129; | |
| mov.b32 %f232, %r127; | |
| mov.b32 %r131, %f217; | |
| mov.b32 %r132, %f561; | |
| div.full.f32 %r130, %r131, %r132; | |
| mov.b32 %f233, %r130; | |
| mov.b32 %r134, %f218; | |
| mov.b32 %r135, %f562; | |
| div.full.f32 %r133, %r134, %r135; | |
| mov.b32 %f234, %r133; | |
| mov.b32 %r137, %f219; | |
| mov.b32 %r138, %f563; | |
| div.full.f32 %r136, %r137, %r138; | |
| mov.b32 %f235, %r136; | |
| mov.b32 %r140, %f220; | |
| mov.b32 %r141, %f564; | |
| div.full.f32 %r139, %r140, %r141; | |
| mov.b32 %f236, %r139; | |
| mov.b32 %r143, %f221; | |
| mov.b32 %r144, %f565; | |
| div.full.f32 %r142, %r143, %r144; | |
| mov.b32 %f237, %r142; | |
| .loc 2 98 22 | |
| add.f32 %f598, %f598, %f222; | |
| add.f32 %f599, %f599, %f223; | |
| add.f32 %f600, %f600, %f224; | |
| add.f32 %f601, %f601, %f225; | |
| add.f32 %f602, %f602, %f226; | |
| add.f32 %f603, %f603, %f227; | |
| add.f32 %f604, %f604, %f228; | |
| add.f32 %f605, %f605, %f229; | |
| add.f32 %f606, %f606, %f230; | |
| add.f32 %f607, %f607, %f231; | |
| add.f32 %f608, %f608, %f232; | |
| add.f32 %f609, %f609, %f233; | |
| add.f32 %f610, %f610, %f234; | |
| add.f32 %f611, %f611, %f235; | |
| add.f32 %f612, %f612, %f236; | |
| add.f32 %f613, %f613, %f237; | |
| .loc 2 101 30 | |
| sub.f32 %f238, %f190, %f598; | |
| sub.f32 %f239, %f191, %f599; | |
| sub.f32 %f240, %f192, %f600; | |
| sub.f32 %f241, %f193, %f601; | |
| sub.f32 %f242, %f194, %f602; | |
| sub.f32 %f243, %f195, %f603; | |
| sub.f32 %f244, %f196, %f604; | |
| sub.f32 %f245, %f197, %f605; | |
| sub.f32 %f246, %f198, %f606; | |
| sub.f32 %f247, %f199, %f607; | |
| sub.f32 %f248, %f200, %f608; | |
| sub.f32 %f249, %f201, %f609; | |
| sub.f32 %f250, %f202, %f610; | |
| sub.f32 %f251, %f203, %f611; | |
| sub.f32 %f252, %f204, %f612; | |
| sub.f32 %f253, %f205, %f613; | |
| $L__tmp2: | |
| .loc 1 47 48 | |
| fma.rn.f32 %f582, %f206, %f238, %f582; | |
| fma.rn.f32 %f583, %f207, %f239, %f583; | |
| fma.rn.f32 %f584, %f208, %f240, %f584; | |
| fma.rn.f32 %f585, %f209, %f241, %f585; | |
| fma.rn.f32 %f586, %f210, %f242, %f586; | |
| fma.rn.f32 %f587, %f211, %f243, %f587; | |
| fma.rn.f32 %f588, %f212, %f244, %f588; | |
| fma.rn.f32 %f589, %f213, %f245, %f589; | |
| fma.rn.f32 %f590, %f214, %f246, %f590; | |
| fma.rn.f32 %f591, %f215, %f247, %f591; | |
| fma.rn.f32 %f592, %f216, %f248, %f592; | |
| fma.rn.f32 %f593, %f217, %f249, %f593; | |
| fma.rn.f32 %f594, %f218, %f250, %f594; | |
| fma.rn.f32 %f595, %f219, %f251, %f595; | |
| fma.rn.f32 %f596, %f220, %f252, %f596; | |
| fma.rn.f32 %f597, %f221, %f253, %f597; | |
| .loc 1 31 36 | |
| add.s64 %rd105, %rd105, 256; | |
| add.s32 %r406, %r406, 64; | |
| setp.lt.u32 %p62, %r406, 192; | |
| @%p62 bra $L__BB0_1; | |
| bra.uni $L__BB0_4; | |
| $L__BB0_1: | |
| .loc 1 39 40 | |
| setp.lt.u64 %p41, %rd1, 50257; | |
| .loc 1 35 34 | |
| add.s64 %rd67, %rd6, %rd105; | |
| add.s64 %rd68, %rd67, 16; | |
| add.s64 %rd69, %rd5, %rd105; | |
| .loc 1 35 50 | |
| add.s64 %rd70, %rd69, 16; | |
| mov.b32 %r342, 0; | |
| mov.u32 %r32, 0x0; | |
| mov.u32 %r33, 0x0; | |
| mov.u32 %r34, 0x0; | |
| mov.u32 %r35, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r32, %r33, %r34, %r35 }, [ %rd67 + 0 ]; | |
| @!%p1 mov.u32 %r32, %r342; | |
| @!%p1 mov.u32 %r33, %r342; | |
| @!%p1 mov.u32 %r34, %r342; | |
| @!%p1 mov.u32 %r35, %r342; | |
| mov.b32 %f65, %r32; | |
| mov.b32 %f66, %r33; | |
| mov.b32 %f67, %r34; | |
| mov.b32 %f68, %r35; | |
| mov.u32 %r40, 0x0; | |
| mov.u32 %r41, 0x0; | |
| mov.u32 %r42, 0x0; | |
| mov.u32 %r43, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r40, %r41, %r42, %r43 }, [ %rd68 + 0 ]; | |
| @!%p1 mov.u32 %r40, %r342; | |
| @!%p1 mov.u32 %r41, %r342; | |
| @!%p1 mov.u32 %r42, %r342; | |
| @!%p1 mov.u32 %r43, %r342; | |
| mov.b32 %f69, %r40; | |
| mov.b32 %f70, %r41; | |
| mov.b32 %f71, %r42; | |
| mov.b32 %f72, %r43; | |
| mov.u32 %r48, 0x0; | |
| mov.u32 %r49, 0x0; | |
| mov.u32 %r50, 0x0; | |
| mov.u32 %r51, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r48, %r49, %r50, %r51 }, [ %rd69 + 0 ]; | |
| @!%p1 mov.u32 %r48, %r342; | |
| @!%p1 mov.u32 %r49, %r342; | |
| @!%p1 mov.u32 %r50, %r342; | |
| @!%p1 mov.u32 %r51, %r342; | |
| mov.b32 %f73, %r48; | |
| mov.b32 %f74, %r49; | |
| mov.b32 %f75, %r50; | |
| mov.b32 %f76, %r51; | |
| mov.u32 %r56, 0x0; | |
| mov.u32 %r57, 0x0; | |
| mov.u32 %r58, 0x0; | |
| mov.u32 %r59, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r56, %r57, %r58, %r59 }, [ %rd70 + 0 ]; | |
| @!%p1 mov.u32 %r56, %r342; | |
| @!%p1 mov.u32 %r57, %r342; | |
| @!%p1 mov.u32 %r58, %r342; | |
| @!%p1 mov.u32 %r59, %r342; | |
| mov.b32 %f77, %r56; | |
| mov.b32 %f78, %r57; | |
| mov.b32 %f79, %r58; | |
| mov.b32 %f80, %r59; | |
| mov.b32 %r405, 883; | |
| mov.u64 %rd104, 1; | |
| .loc 1 39 55 | |
| @%p41 bra $L__BB0_3; | |
| mov.u64 %rd71, assertMessage_0; | |
| cvta.global.u64 %rd72, %rd71; | |
| mov.u64 %rd73, assertFile_0; | |
| cvta.global.u64 %rd74, %rd73; | |
| mov.u64 %rd75, assertFunc_0; | |
| cvta.global.u64 %rd76, %rd75; | |
| { // callseq 6, 0 | |
| .reg .b32 temp_param_reg; | |
| .param .b64 param0; | |
| st.param.b64 [param0+0], %rd72; | |
| .param .b64 param1; | |
| st.param.b64 [param1+0], %rd74; | |
| .param .b32 param2; | |
| st.param.b32 [param2+0], %r405; | |
| .param .b64 param3; | |
| st.param.b64 [param3+0], %rd76; | |
| .param .b64 param4; | |
| st.param.b64 [param4+0], %rd104; | |
| call.uni | |
| __assertfail, | |
| ( | |
| param0, | |
| param1, | |
| param2, | |
| param3, | |
| param4 | |
| ); | |
| } // callseq 6 | |
| bra.uni $L__BB0_3; | |
| $L__BB0_4: | |
| .loc 1 31 36 | |
| and.b32 %r254, %r4, 3; | |
| mad.lo.s32 %r255, %r254, 72, %r2; | |
| shl.b32 %r256, %r255, 2; | |
| mov.u32 %r257, global_smem; | |
| add.s32 %r258, %r257, %r256; | |
| st.shared.f32 [%r258], %f566; | |
| st.shared.f32 [%r258+1152], %f567; | |
| st.shared.f32 [%r258+2304], %f568; | |
| st.shared.f32 [%r258+3456], %f569; | |
| st.shared.f32 [%r258+4608], %f570; | |
| st.shared.f32 [%r258+5760], %f571; | |
| st.shared.f32 [%r258+6912], %f572; | |
| st.shared.f32 [%r258+8064], %f573; | |
| bar.sync 0; | |
| mad.lo.s32 %r259, %r1, 72, %r3; | |
| shl.b32 %r260, %r259, 2; | |
| add.s32 %r261, %r257, %r260; | |
| ld.shared.v4.f32 {%f254, %f255, %f256, %f257}, [%r261]; | |
| ld.shared.v4.f32 {%f258, %f259, %f260, %f261}, [%r261+16]; | |
| bar.sync 0; | |
| st.shared.f32 [%r258], %f574; | |
| st.shared.f32 [%r258+1152], %f575; | |
| st.shared.f32 [%r258+2304], %f576; | |
| st.shared.f32 [%r258+3456], %f577; | |
| st.shared.f32 [%r258+4608], %f578; | |
| st.shared.f32 [%r258+5760], %f579; | |
| st.shared.f32 [%r258+6912], %f580; | |
| st.shared.f32 [%r258+8064], %f581; | |
| bar.sync 0; | |
| ld.shared.v4.f32 {%f262, %f263, %f264, %f265}, [%r261]; | |
| ld.shared.v4.f32 {%f266, %f267, %f268, %f269}, [%r261+16]; | |
| $L__tmp3: | |
| .loc 2 108 21 | |
| sub.f32 %f270, %f599, %f598; | |
| .loc 2 109 28 | |
| add.f32 %f271, %f254, %f255; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p63, %f271, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r146, %f255; | |
| mov.b32 %r147, %f271; | |
| div.full.f32 %r145, %r146, %r147; | |
| mov.b32 %f272, %r145; | |
| .loc 2 110 49 | |
| selp.f32 %f273, 0f00000000, %f272, %p63; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f274, %f270, %f273, %f598; | |
| .loc 2 113 15 | |
| add.f32 %f275, %f582, %f583; | |
| .loc 2 113 30 | |
| mul.f32 %f276, %f270, %f270; | |
| .loc 2 113 38 | |
| mul.f32 %f277, %f276, %f254; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f278, %f277, %f273, %f275; | |
| .loc 2 108 21 | |
| sub.f32 %f279, %f600, %f274; | |
| .loc 2 109 28 | |
| add.f32 %f280, %f256, %f271; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p64, %f280, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r150, %f280; | |
| mov.b32 %r149, %f256; | |
| div.full.f32 %r148, %r149, %r150; | |
| mov.b32 %f281, %r148; | |
| .loc 2 110 49 | |
| selp.f32 %f282, 0f00000000, %f281, %p64; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f283, %f282, %f279, %f274; | |
| .loc 2 113 15 | |
| add.f32 %f284, %f584, %f278; | |
| .loc 2 113 30 | |
| mul.f32 %f285, %f279, %f279; | |
| .loc 2 113 38 | |
| mul.f32 %f286, %f271, %f285; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f287, %f282, %f286, %f284; | |
| .loc 2 108 21 | |
| sub.f32 %f288, %f601, %f283; | |
| .loc 2 109 28 | |
| add.f32 %f289, %f257, %f280; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p65, %f289, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r153, %f289; | |
| mov.b32 %r152, %f257; | |
| div.full.f32 %r151, %r152, %r153; | |
| mov.b32 %f290, %r151; | |
| .loc 2 110 49 | |
| selp.f32 %f291, 0f00000000, %f290, %p65; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f292, %f291, %f288, %f283; | |
| .loc 2 113 15 | |
| add.f32 %f293, %f585, %f287; | |
| .loc 2 113 30 | |
| mul.f32 %f294, %f288, %f288; | |
| .loc 2 113 38 | |
| mul.f32 %f295, %f280, %f294; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f296, %f291, %f295, %f293; | |
| .loc 2 108 21 | |
| sub.f32 %f297, %f602, %f292; | |
| .loc 2 109 28 | |
| add.f32 %f298, %f258, %f289; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p66, %f298, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r156, %f298; | |
| mov.b32 %r155, %f258; | |
| div.full.f32 %r154, %r155, %r156; | |
| mov.b32 %f299, %r154; | |
| .loc 2 110 49 | |
| selp.f32 %f300, 0f00000000, %f299, %p66; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f301, %f300, %f297, %f292; | |
| .loc 2 113 15 | |
| add.f32 %f302, %f586, %f296; | |
| .loc 2 113 30 | |
| mul.f32 %f303, %f297, %f297; | |
| .loc 2 113 38 | |
| mul.f32 %f304, %f289, %f303; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f305, %f300, %f304, %f302; | |
| .loc 2 108 21 | |
| sub.f32 %f306, %f603, %f301; | |
| .loc 2 109 28 | |
| add.f32 %f307, %f259, %f298; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p67, %f307, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r159, %f307; | |
| mov.b32 %r158, %f259; | |
| div.full.f32 %r157, %r158, %r159; | |
| mov.b32 %f308, %r157; | |
| .loc 2 110 49 | |
| selp.f32 %f309, 0f00000000, %f308, %p67; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f310, %f309, %f306, %f301; | |
| .loc 2 113 15 | |
| add.f32 %f311, %f587, %f305; | |
| .loc 2 113 30 | |
| mul.f32 %f312, %f306, %f306; | |
| .loc 2 113 38 | |
| mul.f32 %f313, %f298, %f312; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f314, %f309, %f313, %f311; | |
| .loc 2 108 21 | |
| sub.f32 %f315, %f604, %f310; | |
| .loc 2 109 28 | |
| add.f32 %f316, %f260, %f307; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p68, %f316, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r162, %f316; | |
| mov.b32 %r161, %f260; | |
| div.full.f32 %r160, %r161, %r162; | |
| mov.b32 %f317, %r160; | |
| .loc 2 110 49 | |
| selp.f32 %f318, 0f00000000, %f317, %p68; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f319, %f318, %f315, %f310; | |
| .loc 2 113 15 | |
| add.f32 %f320, %f588, %f314; | |
| .loc 2 113 30 | |
| mul.f32 %f321, %f315, %f315; | |
| .loc 2 113 38 | |
| mul.f32 %f322, %f307, %f321; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f323, %f318, %f322, %f320; | |
| .loc 2 108 21 | |
| sub.f32 %f324, %f605, %f319; | |
| .loc 2 109 28 | |
| add.f32 %f325, %f261, %f316; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p69, %f325, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r165, %f325; | |
| mov.b32 %r164, %f261; | |
| div.full.f32 %r163, %r164, %r165; | |
| mov.b32 %f326, %r163; | |
| .loc 2 110 49 | |
| selp.f32 %f327, 0f00000000, %f326, %p69; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f328, %f327, %f324, %f319; | |
| .loc 2 113 15 | |
| add.f32 %f329, %f589, %f323; | |
| .loc 2 113 30 | |
| mul.f32 %f330, %f324, %f324; | |
| .loc 2 113 38 | |
| mul.f32 %f331, %f316, %f330; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f332, %f327, %f331, %f329; | |
| .loc 2 108 21 | |
| sub.f32 %f333, %f607, %f606; | |
| .loc 2 109 28 | |
| add.f32 %f334, %f262, %f263; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p70, %f334, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r167, %f263; | |
| mov.b32 %r168, %f334; | |
| div.full.f32 %r166, %r167, %r168; | |
| mov.b32 %f335, %r166; | |
| .loc 2 110 49 | |
| selp.f32 %f336, 0f00000000, %f335, %p70; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f337, %f333, %f336, %f606; | |
| .loc 2 113 15 | |
| add.f32 %f338, %f590, %f591; | |
| .loc 2 113 30 | |
| mul.f32 %f339, %f333, %f333; | |
| .loc 2 113 38 | |
| mul.f32 %f340, %f339, %f262; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f341, %f340, %f336, %f338; | |
| .loc 2 108 21 | |
| sub.f32 %f342, %f608, %f337; | |
| .loc 2 109 28 | |
| add.f32 %f343, %f264, %f334; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p71, %f343, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r171, %f343; | |
| mov.b32 %r170, %f264; | |
| div.full.f32 %r169, %r170, %r171; | |
| mov.b32 %f344, %r169; | |
| .loc 2 110 49 | |
| selp.f32 %f345, 0f00000000, %f344, %p71; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f346, %f345, %f342, %f337; | |
| .loc 2 113 15 | |
| add.f32 %f347, %f592, %f341; | |
| .loc 2 113 30 | |
| mul.f32 %f348, %f342, %f342; | |
| .loc 2 113 38 | |
| mul.f32 %f349, %f334, %f348; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f350, %f345, %f349, %f347; | |
| .loc 2 108 21 | |
| sub.f32 %f351, %f609, %f346; | |
| .loc 2 109 28 | |
| add.f32 %f352, %f265, %f343; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p72, %f352, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r174, %f352; | |
| mov.b32 %r173, %f265; | |
| div.full.f32 %r172, %r173, %r174; | |
| mov.b32 %f353, %r172; | |
| .loc 2 110 49 | |
| selp.f32 %f354, 0f00000000, %f353, %p72; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f355, %f354, %f351, %f346; | |
| .loc 2 113 15 | |
| add.f32 %f356, %f593, %f350; | |
| .loc 2 113 30 | |
| mul.f32 %f357, %f351, %f351; | |
| .loc 2 113 38 | |
| mul.f32 %f358, %f343, %f357; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f359, %f354, %f358, %f356; | |
| .loc 2 108 21 | |
| sub.f32 %f360, %f610, %f355; | |
| .loc 2 109 28 | |
| add.f32 %f361, %f266, %f352; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p73, %f361, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r177, %f361; | |
| mov.b32 %r176, %f266; | |
| div.full.f32 %r175, %r176, %r177; | |
| mov.b32 %f362, %r175; | |
| .loc 2 110 49 | |
| selp.f32 %f363, 0f00000000, %f362, %p73; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f364, %f363, %f360, %f355; | |
| .loc 2 113 15 | |
| add.f32 %f365, %f594, %f359; | |
| .loc 2 113 30 | |
| mul.f32 %f366, %f360, %f360; | |
| .loc 2 113 38 | |
| mul.f32 %f367, %f352, %f366; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f368, %f363, %f367, %f365; | |
| .loc 2 108 21 | |
| sub.f32 %f369, %f611, %f364; | |
| .loc 2 109 28 | |
| add.f32 %f370, %f267, %f361; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p74, %f370, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r180, %f370; | |
| mov.b32 %r179, %f267; | |
| div.full.f32 %r178, %r179, %r180; | |
| mov.b32 %f371, %r178; | |
| .loc 2 110 49 | |
| selp.f32 %f372, 0f00000000, %f371, %p74; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f373, %f372, %f369, %f364; | |
| .loc 2 113 15 | |
| add.f32 %f374, %f595, %f368; | |
| .loc 2 113 30 | |
| mul.f32 %f375, %f369, %f369; | |
| .loc 2 113 38 | |
| mul.f32 %f376, %f361, %f375; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f377, %f372, %f376, %f374; | |
| .loc 2 108 21 | |
| sub.f32 %f378, %f612, %f373; | |
| .loc 2 109 28 | |
| add.f32 %f379, %f268, %f370; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p75, %f379, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r183, %f379; | |
| mov.b32 %r182, %f268; | |
| div.full.f32 %r181, %r182, %r183; | |
| mov.b32 %f380, %r181; | |
| .loc 2 110 49 | |
| selp.f32 %f381, 0f00000000, %f380, %p75; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f382, %f381, %f378, %f373; | |
| .loc 2 113 15 | |
| add.f32 %f383, %f596, %f377; | |
| .loc 2 113 30 | |
| mul.f32 %f384, %f378, %f378; | |
| .loc 2 113 38 | |
| mul.f32 %f385, %f370, %f384; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f386, %f381, %f385, %f383; | |
| .loc 2 108 21 | |
| sub.f32 %f387, %f613, %f382; | |
| .loc 2 109 28 | |
| add.f32 %f388, %f269, %f379; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p76, %f388, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r186, %f388; | |
| mov.b32 %r185, %f269; | |
| div.full.f32 %r184, %r185, %r186; | |
| mov.b32 %f389, %r184; | |
| .loc 2 110 49 | |
| selp.f32 %f390, 0f00000000, %f389, %p76; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f391, %f390, %f387, %f382; | |
| .loc 2 113 15 | |
| add.f32 %f392, %f597, %f386; | |
| .loc 2 113 30 | |
| mul.f32 %f393, %f387, %f387; | |
| .loc 2 113 38 | |
| mul.f32 %f394, %f379, %f393; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f395, %f390, %f394, %f392; | |
| $L__tmp4: | |
| .loc 2 120 46 | |
| mov.b32 %r262, %f328; | |
| shfl.sync.bfly.b32 %r263, %r262, 4, 31, -1; | |
| mov.b32 %f396, %r263; | |
| mov.b32 %r264, %f332; | |
| shfl.sync.bfly.b32 %r265, %r264, 4, 31, -1; | |
| mov.b32 %f397, %r265; | |
| shfl.sync.bfly.b32 %r188, %r165, 4, 31, -1; | |
| mov.b32 %f398, %r188; | |
| $L__tmp5: | |
| .loc 2 108 21 | |
| sub.f32 %f399, %f396, %f328; | |
| .loc 2 109 28 | |
| add.f32 %f400, %f325, %f398; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p77, %f400, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r189, %f400; | |
| div.full.f32 %r187, %r188, %r189; | |
| mov.b32 %f401, %r187; | |
| .loc 2 110 49 | |
| selp.f32 %f402, 0f00000000, %f401, %p77; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f403, %f402, %f399, %f328; | |
| .loc 2 113 15 | |
| add.f32 %f404, %f332, %f397; | |
| .loc 2 113 30 | |
| mul.f32 %f405, %f399, %f399; | |
| .loc 2 113 38 | |
| mul.f32 %f406, %f325, %f405; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f407, %f402, %f406, %f404; | |
| $L__tmp6: | |
| .loc 2 120 46 | |
| mov.b32 %r266, %f403; | |
| shfl.sync.bfly.b32 %r267, %r266, 2, 31, -1; | |
| mov.b32 %f408, %r267; | |
| mov.b32 %r268, %f407; | |
| shfl.sync.bfly.b32 %r269, %r268, 2, 31, -1; | |
| mov.b32 %f409, %r269; | |
| shfl.sync.bfly.b32 %r191, %r189, 2, 31, -1; | |
| mov.b32 %f410, %r191; | |
| $L__tmp7: | |
| .loc 2 108 21 | |
| sub.f32 %f411, %f408, %f403; | |
| .loc 2 109 28 | |
| add.f32 %f412, %f400, %f410; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p78, %f412, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r192, %f412; | |
| div.full.f32 %r190, %r191, %r192; | |
| mov.b32 %f413, %r190; | |
| .loc 2 110 49 | |
| selp.f32 %f414, 0f00000000, %f413, %p78; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f415, %f414, %f411, %f403; | |
| .loc 2 113 15 | |
| add.f32 %f416, %f407, %f409; | |
| .loc 2 113 30 | |
| mul.f32 %f417, %f411, %f411; | |
| .loc 2 113 38 | |
| mul.f32 %f418, %f400, %f417; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f419, %f414, %f418, %f416; | |
| $L__tmp8: | |
| .loc 2 120 46 | |
| mov.b32 %r270, %f415; | |
| shfl.sync.bfly.b32 %r271, %r270, 1, 31, -1; | |
| mov.b32 %f420, %r271; | |
| mov.b32 %r272, %f419; | |
| shfl.sync.bfly.b32 %r273, %r272, 1, 31, -1; | |
| mov.b32 %f421, %r273; | |
| shfl.sync.bfly.b32 %r194, %r192, 1, 31, -1; | |
| mov.b32 %f422, %r194; | |
| $L__tmp9: | |
| .loc 2 108 21 | |
| sub.f32 %f423, %f420, %f415; | |
| .loc 2 109 28 | |
| add.f32 %f424, %f412, %f422; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p79, %f424, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r195, %f424; | |
| div.full.f32 %r193, %r194, %r195; | |
| mov.b32 %f425, %r193; | |
| .loc 2 110 49 | |
| selp.f32 %f426, 0f00000000, %f425, %p79; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f145, %f423, %f426, %f415; | |
| .loc 2 113 15 | |
| add.f32 %f427, %f419, %f421; | |
| .loc 2 113 30 | |
| mul.f32 %f428, %f423, %f423; | |
| .loc 2 113 38 | |
| mul.f32 %f429, %f412, %f428; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f430, %f426, %f429, %f427; | |
| $L__tmp10: | |
| .loc 2 120 46 | |
| mov.b32 %r274, %f391; | |
| shfl.sync.bfly.b32 %r275, %r274, 4, 31, -1; | |
| mov.b32 %f431, %r275; | |
| mov.b32 %r276, %f395; | |
| shfl.sync.bfly.b32 %r277, %r276, 4, 31, -1; | |
| mov.b32 %f432, %r277; | |
| shfl.sync.bfly.b32 %r197, %r186, 4, 31, -1; | |
| mov.b32 %f433, %r197; | |
| $L__tmp11: | |
| .loc 2 108 21 | |
| sub.f32 %f434, %f431, %f391; | |
| .loc 2 109 28 | |
| add.f32 %f435, %f388, %f433; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p80, %f435, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r198, %f435; | |
| div.full.f32 %r196, %r197, %r198; | |
| mov.b32 %f436, %r196; | |
| .loc 2 110 49 | |
| selp.f32 %f437, 0f00000000, %f436, %p80; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f438, %f434, %f437, %f391; | |
| .loc 2 113 15 | |
| add.f32 %f439, %f395, %f432; | |
| .loc 2 113 30 | |
| mul.f32 %f440, %f434, %f434; | |
| .loc 2 113 38 | |
| mul.f32 %f441, %f388, %f440; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f442, %f441, %f437, %f439; | |
| $L__tmp12: | |
| .loc 2 120 46 | |
| mov.b32 %r278, %f438; | |
| shfl.sync.bfly.b32 %r279, %r278, 2, 31, -1; | |
| mov.b32 %f443, %r279; | |
| mov.b32 %r280, %f442; | |
| shfl.sync.bfly.b32 %r281, %r280, 2, 31, -1; | |
| mov.b32 %f444, %r281; | |
| shfl.sync.bfly.b32 %r200, %r198, 2, 31, -1; | |
| mov.b32 %f445, %r200; | |
| $L__tmp13: | |
| .loc 2 108 21 | |
| sub.f32 %f446, %f443, %f438; | |
| .loc 2 109 28 | |
| add.f32 %f447, %f435, %f445; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p81, %f447, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r201, %f447; | |
| div.full.f32 %r199, %r200, %r201; | |
| mov.b32 %f448, %r199; | |
| .loc 2 110 49 | |
| selp.f32 %f449, 0f00000000, %f448, %p81; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f450, %f446, %f449, %f438; | |
| .loc 2 113 15 | |
| add.f32 %f451, %f442, %f444; | |
| .loc 2 113 30 | |
| mul.f32 %f452, %f446, %f446; | |
| .loc 2 113 38 | |
| mul.f32 %f453, %f435, %f452; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f454, %f449, %f453, %f451; | |
| $L__tmp14: | |
| .loc 2 120 46 | |
| mov.b32 %r282, %f450; | |
| shfl.sync.bfly.b32 %r283, %r282, 1, 31, -1; | |
| mov.b32 %f455, %r283; | |
| mov.b32 %r284, %f454; | |
| shfl.sync.bfly.b32 %r285, %r284, 1, 31, -1; | |
| mov.b32 %f456, %r285; | |
| shfl.sync.bfly.b32 %r203, %r201, 1, 31, -1; | |
| mov.b32 %f457, %r203; | |
| $L__tmp15: | |
| .loc 2 108 21 | |
| sub.f32 %f458, %f455, %f450; | |
| .loc 2 109 28 | |
| add.f32 %f459, %f447, %f457; | |
| .loc 2 110 39 | |
| setp.eq.f32 %p82, %f459, 0f00000000; | |
| .loc 2 110 60 | |
| mov.b32 %r204, %f459; | |
| div.full.f32 %r202, %r203, %r204; | |
| mov.b32 %f460, %r202; | |
| .loc 2 110 49 | |
| selp.f32 %f461, 0f00000000, %f460, %p82; | |
| .loc 2 112 17 | |
| fma.rn.f32 %f146, %f458, %f461, %f450; | |
| .loc 2 113 15 | |
| add.f32 %f462, %f454, %f456; | |
| .loc 2 113 30 | |
| mul.f32 %f463, %f458, %f458; | |
| .loc 2 113 38 | |
| mul.f32 %f464, %f447, %f463; | |
| .loc 2 113 22 | |
| fma.rn.f32 %f465, %f461, %f464, %f462; | |
| $L__tmp16: | |
| .loc 1 69 23 | |
| mov.b32 %r206, %f430; | |
| mov.b32 %r207, 1132462080; | |
| div.full.f32 %r205, %r206, %r207; | |
| mov.b32 %f466, %r205; | |
| mov.b32 %r230, %f465; | |
| div.full.f32 %r229, %r230, %r207; | |
| mov.b32 %f467, %r229; | |
| .loc 1 71 24 | |
| add.f32 %f147, %f466, 0f3727C5AC; | |
| add.f32 %f148, %f467, 0f3727C5AC; | |
| .loc 1 55 36 | |
| add.s64 %rd9, %rd12, %rd2; | |
| shl.b32 %r286, %r11, 14; | |
| shl.b32 %r287, %r1, 8; | |
| or.b32 %r288, %r286, %r287; | |
| or.b32 %r8, %r288, %r3; | |
| mov.u64 %rd106, 0; | |
| mov.b32 %r407, -64; | |
| rsqrt.approx.ftz.f32 %f516, %f147; | |
| rsqrt.approx.ftz.f32 %f517, %f148; | |
| bra.uni $L__BB0_5; | |
| $L__BB0_7: | |
| .loc 1 65 35 | |
| add.s64 %rd96, %rd4, %rd106; | |
| add.s64 %rd97, %rd96, 16; | |
| add.s64 %rd98, %rd3, %rd106; | |
| .loc 1 65 54 | |
| add.s64 %rd99, %rd98, 16; | |
| mov.u32 %r338, 0x0; | |
| mov.u32 %r339, 0x0; | |
| mov.u32 %r340, 0x0; | |
| mov.u32 %r341, 0x0; | |
| @%p1 ld.global.L1::evict_first.v4.b32 { %r338, %r339, %r340, %r341 }, [ %rd96 + 0 ]; | |
| @!%p1 mov.u32 %r338, %r342; | |
| @!%p1 mov.u32 %r339, %r342; | |
| @!%p1 mov.u32 %r340, %r342; | |
| @!%p1 mov.u32 %r341, %r342; | |
| mov.b32 %f468, %r338; | |
| mov.b32 %f469, %r339; | |
| mov.b32 %f470, %r340; | |
| mov.b32 %f471, %r341; | |
| mov.u32 %r346, 0x0; | |
| mov.u32 %r347, 0x0; | |
| mov.u32 %r348, 0x0; | |
| mov.u32 %r349, 0x0; | |
| @%p1 ld.global.L1::evict_first.v4.b32 { %r346, %r347, %r348, %r349 }, [ %rd97 + 0 ]; | |
| @!%p1 mov.u32 %r346, %r342; | |
| @!%p1 mov.u32 %r347, %r342; | |
| @!%p1 mov.u32 %r348, %r342; | |
| @!%p1 mov.u32 %r349, %r342; | |
| mov.b32 %f472, %r346; | |
| mov.b32 %f473, %r347; | |
| mov.b32 %f474, %r348; | |
| mov.b32 %f475, %r349; | |
| mov.u32 %r354, 0x0; | |
| mov.u32 %r355, 0x0; | |
| mov.u32 %r356, 0x0; | |
| mov.u32 %r357, 0x0; | |
| @%p1 ld.global.L1::evict_first.v4.b32 { %r354, %r355, %r356, %r357 }, [ %rd98 + 0 ]; | |
| @!%p1 mov.u32 %r354, %r342; | |
| @!%p1 mov.u32 %r355, %r342; | |
| @!%p1 mov.u32 %r356, %r342; | |
| @!%p1 mov.u32 %r357, %r342; | |
| mov.b32 %f476, %r354; | |
| mov.b32 %f477, %r355; | |
| mov.b32 %f478, %r356; | |
| mov.b32 %f479, %r357; | |
| mov.u32 %r362, 0x0; | |
| mov.u32 %r363, 0x0; | |
| mov.u32 %r364, 0x0; | |
| mov.u32 %r365, 0x0; | |
| @%p1 ld.global.L1::evict_first.v4.b32 { %r362, %r363, %r364, %r365 }, [ %rd99 + 0 ]; | |
| @!%p1 mov.u32 %r362, %r342; | |
| @!%p1 mov.u32 %r363, %r342; | |
| @!%p1 mov.u32 %r364, %r342; | |
| @!%p1 mov.u32 %r365, %r342; | |
| mov.b32 %f480, %r362; | |
| mov.b32 %f481, %r363; | |
| mov.b32 %f482, %r364; | |
| mov.b32 %f483, %r365; | |
| .loc 1 66 24 | |
| add.f32 %f484, %f149, %f468; | |
| add.f32 %f485, %f150, %f469; | |
| add.f32 %f486, %f151, %f470; | |
| add.f32 %f487, %f152, %f471; | |
| add.f32 %f488, %f153, %f472; | |
| add.f32 %f489, %f154, %f473; | |
| add.f32 %f490, %f155, %f474; | |
| add.f32 %f491, %f156, %f475; | |
| add.f32 %f492, %f157, %f476; | |
| add.f32 %f493, %f158, %f477; | |
| add.f32 %f494, %f159, %f478; | |
| add.f32 %f495, %f160, %f479; | |
| add.f32 %f496, %f161, %f480; | |
| add.f32 %f497, %f162, %f481; | |
| add.f32 %f498, %f163, %f482; | |
| add.f32 %f499, %f164, %f483; | |
| .loc 1 67 24 | |
| sub.f32 %f500, %f484, %f145; | |
| sub.f32 %f501, %f485, %f145; | |
| sub.f32 %f502, %f486, %f145; | |
| sub.f32 %f503, %f487, %f145; | |
| sub.f32 %f504, %f488, %f145; | |
| sub.f32 %f505, %f489, %f145; | |
| sub.f32 %f506, %f490, %f145; | |
| sub.f32 %f507, %f491, %f145; | |
| sub.f32 %f508, %f492, %f146; | |
| sub.f32 %f509, %f493, %f146; | |
| sub.f32 %f510, %f494, %f146; | |
| sub.f32 %f511, %f495, %f146; | |
| sub.f32 %f512, %f496, %f146; | |
| sub.f32 %f513, %f497, %f146; | |
| sub.f32 %f514, %f498, %f146; | |
| sub.f32 %f515, %f499, %f146; | |
| .loc 1 73 24 | |
| mul.f32 %f518, %f500, %f516; | |
| mul.f32 %f519, %f501, %f516; | |
| mul.f32 %f520, %f502, %f516; | |
| mul.f32 %f521, %f503, %f516; | |
| mul.f32 %f522, %f504, %f516; | |
| mul.f32 %f523, %f505, %f516; | |
| mul.f32 %f524, %f506, %f516; | |
| mul.f32 %f525, %f507, %f516; | |
| mul.f32 %f526, %f508, %f517; | |
| mul.f32 %f527, %f509, %f517; | |
| mul.f32 %f528, %f510, %f517; | |
| mul.f32 %f529, %f511, %f517; | |
| mul.f32 %f530, %f512, %f517; | |
| mul.f32 %f531, %f513, %f517; | |
| mul.f32 %f532, %f514, %f517; | |
| mul.f32 %f533, %f515, %f517; | |
| .loc 1 74 24 | |
| mul.f32 %f534, %f518, %f165; | |
| mul.f32 %f535, %f519, %f166; | |
| mul.f32 %f536, %f520, %f167; | |
| mul.f32 %f537, %f521, %f168; | |
| mul.f32 %f538, %f522, %f169; | |
| mul.f32 %f539, %f523, %f170; | |
| mul.f32 %f540, %f524, %f171; | |
| mul.f32 %f541, %f525, %f172; | |
| mul.f32 %f542, %f526, %f165; | |
| mul.f32 %f543, %f527, %f166; | |
| mul.f32 %f544, %f528, %f167; | |
| mul.f32 %f545, %f529, %f168; | |
| mul.f32 %f546, %f530, %f169; | |
| mul.f32 %f547, %f531, %f170; | |
| mul.f32 %f548, %f532, %f171; | |
| mul.f32 %f549, %f533, %f172; | |
| .loc 1 76 35 | |
| add.s32 %r394, %r8, %r407; | |
| add.s32 %r395, %r394, 64; | |
| .loc 1 76 29 | |
| add.s32 %r396, %r394, 8256; | |
| mul.wide.s32 %rd102, %r395, 2; | |
| add.s64 %rd100, %rd13, %rd102; | |
| mul.wide.s32 %rd103, %r396, 2; | |
| add.s64 %rd101, %rd13, %rd103; | |
| .loc 1 76 52 | |
| mov.b32 %r370, %f534; | |
| cvt.rn.bf16.f32 %rs1, %r370; | |
| mov.b32 %r371, %f535; | |
| cvt.rn.bf16.f32 %rs2, %r371; | |
| mov.b32 %r372, %f536; | |
| cvt.rn.bf16.f32 %rs3, %r372; | |
| mov.b32 %r373, %f537; | |
| cvt.rn.bf16.f32 %rs4, %r373; | |
| mov.b32 %r374, %f538; | |
| cvt.rn.bf16.f32 %rs5, %r374; | |
| mov.b32 %r375, %f539; | |
| cvt.rn.bf16.f32 %rs6, %r375; | |
| mov.b32 %r376, %f540; | |
| cvt.rn.bf16.f32 %rs7, %r376; | |
| mov.b32 %r377, %f541; | |
| cvt.rn.bf16.f32 %rs8, %r377; | |
| mov.b32 %r378, %f542; | |
| cvt.rn.bf16.f32 %rs9, %r378; | |
| mov.b32 %r379, %f543; | |
| cvt.rn.bf16.f32 %rs10, %r379; | |
| mov.b32 %r380, %f544; | |
| cvt.rn.bf16.f32 %rs11, %r380; | |
| mov.b32 %r381, %f545; | |
| cvt.rn.bf16.f32 %rs12, %r381; | |
| mov.b32 %r382, %f546; | |
| cvt.rn.bf16.f32 %rs13, %r382; | |
| mov.b32 %r383, %f547; | |
| cvt.rn.bf16.f32 %rs14, %r383; | |
| mov.b32 %r384, %f548; | |
| cvt.rn.bf16.f32 %rs15, %r384; | |
| mov.b32 %r385, %f549; | |
| cvt.rn.bf16.f32 %rs16, %r385; | |
| mov.b32 %r397, {%rs1, %rs2}; | |
| mov.b32 %r398, {%rs3, %rs4}; | |
| mov.b32 %r399, {%rs5, %rs6}; | |
| mov.b32 %r400, {%rs7, %rs8}; | |
| @%p1 st.global.v4.b32 [ %rd100 + 0 ], { %r397, %r398, %r399, %r400 }; | |
| mov.b32 %r401, {%rs9, %rs10}; | |
| mov.b32 %r402, {%rs11, %rs12}; | |
| mov.b32 %r403, {%rs13, %rs14}; | |
| mov.b32 %r404, {%rs15, %rs16}; | |
| @%p1 st.global.v4.b32 [ %rd101 + 0 ], { %r401, %r402, %r403, %r404 }; | |
| .loc 1 55 36 | |
| add.s64 %rd106, %rd106, 256; | |
| add.s32 %r407, %r407, 64; | |
| setp.lt.u32 %p136, %r407, 192; | |
| @%p136 bra $L__BB0_5; | |
| bra.uni $L__BB0_8; | |
| $L__BB0_5: | |
| .loc 1 59 35 | |
| add.s64 %rd83, %rd6, %rd106; | |
| add.s64 %rd84, %rd83, 16; | |
| add.s64 %rd85, %rd5, %rd106; | |
| .loc 1 59 51 | |
| add.s64 %rd86, %rd85, 16; | |
| mov.u32 %r289, 0x0; | |
| mov.u32 %r290, 0x0; | |
| mov.u32 %r291, 0x0; | |
| mov.u32 %r292, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r289, %r290, %r291, %r292 }, [ %rd83 + 0 ]; | |
| @!%p1 mov.u32 %r289, %r342; | |
| @!%p1 mov.u32 %r290, %r342; | |
| @!%p1 mov.u32 %r291, %r342; | |
| @!%p1 mov.u32 %r292, %r342; | |
| mov.b32 %f149, %r289; | |
| mov.b32 %f150, %r290; | |
| mov.b32 %f151, %r291; | |
| mov.b32 %f152, %r292; | |
| mov.u32 %r297, 0x0; | |
| mov.u32 %r298, 0x0; | |
| mov.u32 %r299, 0x0; | |
| mov.u32 %r300, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r297, %r298, %r299, %r300 }, [ %rd84 + 0 ]; | |
| @!%p1 mov.u32 %r297, %r342; | |
| @!%p1 mov.u32 %r298, %r342; | |
| @!%p1 mov.u32 %r299, %r342; | |
| @!%p1 mov.u32 %r300, %r342; | |
| mov.b32 %f153, %r297; | |
| mov.b32 %f154, %r298; | |
| mov.b32 %f155, %r299; | |
| mov.b32 %f156, %r300; | |
| mov.u32 %r305, 0x0; | |
| mov.u32 %r306, 0x0; | |
| mov.u32 %r307, 0x0; | |
| mov.u32 %r308, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r305, %r306, %r307, %r308 }, [ %rd85 + 0 ]; | |
| @!%p1 mov.u32 %r305, %r342; | |
| @!%p1 mov.u32 %r306, %r342; | |
| @!%p1 mov.u32 %r307, %r342; | |
| @!%p1 mov.u32 %r308, %r342; | |
| mov.b32 %f157, %r305; | |
| mov.b32 %f158, %r306; | |
| mov.b32 %f159, %r307; | |
| mov.b32 %f160, %r308; | |
| mov.u32 %r313, 0x0; | |
| mov.u32 %r314, 0x0; | |
| mov.u32 %r315, 0x0; | |
| mov.u32 %r316, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r313, %r314, %r315, %r316 }, [ %rd86 + 0 ]; | |
| @!%p1 mov.u32 %r313, %r342; | |
| @!%p1 mov.u32 %r314, %r342; | |
| @!%p1 mov.u32 %r315, %r342; | |
| @!%p1 mov.u32 %r316, %r342; | |
| mov.b32 %f161, %r313; | |
| mov.b32 %f162, %r314; | |
| mov.b32 %f163, %r315; | |
| mov.b32 %f164, %r316; | |
| .loc 1 60 35 | |
| add.s64 %rd87, %rd9, %rd106; | |
| .loc 1 60 40 | |
| add.s64 %rd88, %rd87, 16; | |
| mov.u32 %r321, 0x0; | |
| mov.u32 %r322, 0x0; | |
| mov.u32 %r323, 0x0; | |
| mov.u32 %r324, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r321, %r322, %r323, %r324 }, [ %rd87 + 0 ]; | |
| @!%p1 mov.u32 %r321, %r342; | |
| @!%p1 mov.u32 %r322, %r342; | |
| @!%p1 mov.u32 %r323, %r342; | |
| @!%p1 mov.u32 %r324, %r342; | |
| mov.b32 %f165, %r321; | |
| mov.b32 %f166, %r322; | |
| mov.b32 %f167, %r323; | |
| mov.b32 %f168, %r324; | |
| mov.u32 %r329, 0x0; | |
| mov.u32 %r330, 0x0; | |
| mov.u32 %r331, 0x0; | |
| mov.u32 %r332, 0x0; | |
| @%p1 ld.global.L1::evict_last.v4.b32 { %r329, %r330, %r331, %r332 }, [ %rd88 + 0 ]; | |
| @!%p1 mov.u32 %r329, %r342; | |
| @!%p1 mov.u32 %r330, %r342; | |
| @!%p1 mov.u32 %r331, %r342; | |
| @!%p1 mov.u32 %r332, %r342; | |
| mov.b32 %f169, %r329; | |
| mov.b32 %f170, %r330; | |
| mov.b32 %f171, %r331; | |
| mov.b32 %f172, %r332; | |
| .loc 1 64 57 | |
| @%p41 bra $L__BB0_7; | |
| mov.u64 %rd89, assertMessage_1; | |
| cvta.global.u64 %rd90, %rd89; | |
| mov.u64 %rd91, assertFile_1; | |
| cvta.global.u64 %rd92, %rd91; | |
| mov.u64 %rd93, assertFunc_1; | |
| cvta.global.u64 %rd94, %rd93; | |
| { // callseq 7, 0 | |
| .reg .b32 temp_param_reg; | |
| .param .b64 param0; | |
| st.param.b64 [param0+0], %rd90; | |
| .param .b64 param1; | |
| st.param.b64 [param1+0], %rd92; | |
| .param .b32 param2; | |
| st.param.b32 [param2+0], %r405; | |
| .param .b64 param3; | |
| st.param.b64 [param3+0], %rd94; | |
| .param .b64 param4; | |
| st.param.b64 [param4+0], %rd104; | |
| call.uni | |
| __assertfail, | |
| ( | |
| param0, | |
| param1, | |
| param2, | |
| param3, | |
| param4 | |
| ); | |
| } // callseq 7 | |
| bra.uni $L__BB0_7; | |
| $L__BB0_8: | |
| .loc 1 55 4 | |
| ret; | |
| $L__tmp17: | |
| $L__func_end0: | |
| } | |
| // .globl __nv_rsqrtf | |
| .visible .func (.param .b32 func_retval0) __nv_rsqrtf( | |
| .param .b32 __nv_rsqrtf_param_0 | |
| ) | |
| { | |
| .reg .f32 %f<3>; | |
| $L__func_begin1: | |
| ld.param.f32 %f1, [__nv_rsqrtf_param_0]; | |
| rsqrt.approx.ftz.f32 %f2, %f1; | |
| st.param.f32 [func_retval0+0], %f2; | |
| ret; | |
| $L__func_end1: | |
| } | |
| .file 1 "/tmp/torchinductor_root/gx/cgx5lxpuexpindj4dsmjz5x42uhyy7iskevq7ovzpwagb3t5powj.py" | |
| .file 2 "/usr/local/lib/python3.10/dist-packages/torch/_inductor/triton_helpers.py" | |
| .section .debug_abbrev | |
| { | |
| .b8 1 | |
| .b8 17 | |
| .b8 1 | |
| .b8 37 | |
| .b8 8 | |
| .b8 19 | |
| .b8 5 | |
| .b8 3 | |
| .b8 8 | |
| .b8 16 | |
| .b8 6 | |
| .b8 27 | |
| .b8 8 | |
| .b8 180 | |
| .b8 66 | |
| .b8 12 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 0 | |
| .b8 0 | |
| .b8 2 | |
| .b8 46 | |
| .b8 0 | |
| .b8 135 | |
| .b8 64 | |
| .b8 8 | |
| .b8 3 | |
| .b8 8 | |
| .b8 58 | |
| .b8 11 | |
| .b8 59 | |
| .b8 11 | |
| .b8 63 | |
| .b8 12 | |
| .b8 32 | |
| .b8 11 | |
| .b8 0 | |
| .b8 0 | |
| .b8 3 | |
| .b8 46 | |
| .b8 1 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 64 | |
| .b8 10 | |
| .b8 49 | |
| .b8 19 | |
| .b8 0 | |
| .b8 0 | |
| .b8 4 | |
| .b8 29 | |
| .b8 0 | |
| .b8 49 | |
| .b8 19 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 88 | |
| .b8 11 | |
| .b8 89 | |
| .b8 11 | |
| .b8 87 | |
| .b8 11 | |
| .b8 0 | |
| .b8 0 | |
| .b8 5 | |
| .b8 29 | |
| .b8 1 | |
| .b8 49 | |
| .b8 19 | |
| .b8 17 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 88 | |
| .b8 11 | |
| .b8 89 | |
| .b8 11 | |
| .b8 87 | |
| .b8 11 | |
| .b8 0 | |
| .b8 0 | |
| .b8 0 | |
| } | |
| .section .debug_info | |
| { | |
| .b32 298 | |
| .b8 2 | |
| .b8 0 | |
| .b32 .debug_abbrev | |
| .b8 8 | |
| .b8 1 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 0 | |
| .b8 2 | |
| .b8 0 | |
| .b8 99 | |
| .b8 103 | |
| .b8 120 | |
| .b8 53 | |
| .b8 108 | |
| .b8 120 | |
| .b8 112 | |
| .b8 117 | |
| .b8 101 | |
| .b8 120 | |
| .b8 112 | |
| .b8 105 | |
| .b8 110 | |
| .b8 100 | |
| .b8 106 | |
| .b8 52 | |
| .b8 100 | |
| .b8 115 | |
| .b8 109 | |
| .b8 106 | |
| .b8 122 | |
| .b8 53 | |
| .b8 120 | |
| .b8 52 | |
| .b8 50 | |
| .b8 117 | |
| .b8 104 | |
| .b8 121 | |
| .b8 121 | |
| .b8 55 | |
| .b8 105 | |
| .b8 115 | |
| .b8 107 | |
| .b8 101 | |
| .b8 118 | |
| .b8 113 | |
| .b8 55 | |
| .b8 111 | |
| .b8 118 | |
| .b8 122 | |
| .b8 112 | |
| .b8 119 | |
| .b8 97 | |
| .b8 103 | |
| .b8 98 | |
| .b8 51 | |
| .b8 116 | |
| .b8 53 | |
| .b8 112 | |
| .b8 111 | |
| .b8 119 | |
| .b8 106 | |
| .b8 46 | |
| .b8 112 | |
| .b8 121 | |
| .b8 0 | |
| .b32 .debug_line | |
| .b8 47 | |
| .b8 116 | |
| .b8 109 | |
| .b8 112 | |
| .b8 47 | |
| .b8 116 | |
| .b8 111 | |
| .b8 114 | |
| .b8 99 | |
| .b8 104 | |
| .b8 105 | |
| .b8 110 | |
| .b8 100 | |
| .b8 117 | |
| .b8 99 | |
| .b8 116 | |
| .b8 111 | |
| .b8 114 | |
| .b8 95 | |
| .b8 114 | |
| .b8 111 | |
| .b8 111 | |
| .b8 116 | |
| .b8 47 | |
| .b8 103 | |
| .b8 120 | |
| .b8 0 | |
| .b8 1 | |
| .b64 $L__func_begin0 | |
| .b64 $L__func_end0 | |
| .b8 2 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 95 | |
| .b8 95 | |
| .b8 48 | |
| .b8 100 | |
| .b8 49 | |
| .b8 100 | |
| .b8 50 | |
| .b8 100 | |
| .b8 51 | |
| .b8 100 | |
| .b8 52 | |
| .b8 100 | |
| .b8 53 | |
| .b8 100 | |
| .b8 101 | |
| .b8 54 | |
| .b8 100 | |
| .b8 101 | |
| .b8 0 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 95 | |
| .b8 95 | |
| .b8 48 | |
| .b8 100 | |
| .b8 49 | |
| .b8 100 | |
| .b8 50 | |
| .b8 100 | |
| .b8 51 | |
| .b8 100 | |
| .b8 52 | |
| .b8 100 | |
| .b8 53 | |
| .b8 100 | |
| .b8 101 | |
| .b8 54 | |
| .b8 100 | |
| .b8 101 | |
| .b8 0 | |
| .b8 1 | |
| .b8 18 | |
| .b8 1 | |
| .b8 1 | |
| .b8 3 | |
| .b64 $L__func_begin0 | |
| .b64 $L__func_end0 | |
| .b8 1 | |
| .b8 156 | |
| .b32 125 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp1 | |
| .b64 $L__tmp2 | |
| .b8 2 | |
| .b8 44 | |
| .b8 38 | |
| .b8 5 | |
| .b32 125 | |
| .b64 $L__tmp3 | |
| .b64 $L__tmp16 | |
| .b8 2 | |
| .b8 50 | |
| .b8 41 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp3 | |
| .b64 $L__tmp16 | |
| .b8 2 | |
| .b8 120 | |
| .b8 46 | |
| .b8 0 | |
| .b8 4 | |
| .b32 125 | |
| .b64 $L__tmp4 | |
| .b64 $L__tmp15 | |
| .b8 2 | |
| .b8 50 | |
| .b8 41 | |
| .b8 0 | |
| .b8 0 | |
| } | |
| .section .debug_pubnames | |
| { | |
| .b32 $L__pubNames_end0-$L__pubNames_start0 | |
| $L__pubNames_start0: | |
| .b8 2 | |
| .b8 0 | |
| .b32 .debug_info | |
| .b32 302 | |
| .b32 125 | |
| .b8 116 | |
| .b8 114 | |
| .b8 105 | |
| .b8 116 | |
| .b8 111 | |
| .b8 110 | |
| .b8 95 | |
| .b8 95 | |
| .b8 48 | |
| .b8 100 | |
| .b8 49 | |
| .b8 100 | |
| .b8 50 | |
| .b8 100 | |
| .b8 51 | |
| .b8 100 | |
| .b8 52 | |
| .b8 100 | |
| .b8 53 | |
| .b8 100 | |
| .b8 101 | |
| .b8 54 | |
| .b8 100 | |
| .b8 101 | |
| .b8 0 | |
| .b32 0 | |
| $L__pubNames_end0: | |
| } | |
| .section .debug_pubtypes | |
| { | |
| .b32 $L__pubTypes_end0-$L__pubTypes_start0 | |
| $L__pubTypes_start0: | |
| .b8 2 | |
| .b8 0 | |
| .b32 .debug_info | |
| .b32 302 | |
| .b32 0 | |
| $L__pubTypes_end0: | |
| } | |
| .section .debug_loc { } | |