INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_2E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_3E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_4E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_5E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_6E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_7E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_8E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders2_9E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_872a33146thrust12placeholders3_10E[1]; .visible .entry _Z28dequant_gemv_group128_batch823DequantGemvKernelParams( .param .align 8 .b8 _Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0[80] ) { .reg .pred %p<33>; .reg .b16 %rs<788>; .reg .f32 %f<1522>; .reg .b32 %r<763>; .reg .b64 %rd<110>; // demoted variable .shared .align 16 .b8 _ZZ9gemv_int4ILi4ELi128ELi8EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage[5120]; ld.param.v2.u32 {%r61, %r62}, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+48]; ld.param.v2.u32 {%r63, %r64}, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+56]; ld.param.v2.f32 {%f95, %f96}, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+64]; ld.param.v4.u8 {%rs73, %rs74, %rs75, %rs76}, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+72]; ld.param.u64 %rd36, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+40]; ld.param.u64 %rd35, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+32]; ld.param.u64 %rd34, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+24]; ld.param.u64 %rd33, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+16]; ld.param.u64 %rd32, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+8]; cvta.to.global.u64 %rd2, %rd32; mov.u32 %r762, %tid.y; shl.b32 %r65, %r762, 5; mov.u32 %r66, %tid.x; add.s32 %r761, %r65, %r66; shl.b32 %r760, %r761, 2; setp.ge.u32 %p1, %r760, %r63; mov.f32 %f1490, 0f00000000; mov.f32 %f1491, %f1490; mov.f32 %f1492, %f1490; mov.f32 %f1493, %f1490; mov.f32 %f1494, %f1490; mov.f32 %f1495, %f1490; mov.f32 %f1496, %f1490; mov.f32 %f1497, %f1490; @%p1 bra $L__BB0_9; cvta.to.global.u64 %rd3, %rd33; mov.u32 %r67, %ctaid.x; mul.lo.s32 %r6, %r64, %r67; $L__BB0_2: mad.lo.s32 %r73, %r63, %r67, %r760; mul.wide.u32 %rd43, %r73, 4; add.s64 %rd38, %rd34, %rd43; // begin inline asm createpolicy.fractional.L2::evict_first.L2::evict_first.b64 %rd37, 1.0; // end inline asm // begin inline asm ld.global.L1::no_allocate.L2::cache_hint.v4.u32 {%r68,%r69,%r70,%r71}, [%rd38], %rd37; // end inline asm shr.u32 %r75, %r66, 2; shl.b32 %r76, %r762, 3; add.s32 %r14, %r76, %r75; add.s32 %r15, %r14, %r6; mul.wide.s32 %rd44, %r15, 2; add.s64 %rd41, %rd36, %rd44; // begin inline asm createpolicy.fractional.L2::evict_first.L2::evict_first.b64 %rd40, 1.0; // end inline asm // begin inline asm ld.global.L1::evict_last.L2::cache_hint.u16 %rs81, [%rd41], %rd40; // end inline asm // begin inline asm { cvt.f32.f16 %f105, %rs81;} // end inline asm shl.b16 %rs787, %rs73, 3; setp.eq.s64 %p2, %rd35, 0; @%p2 bra $L__BB0_4; shr.u32 %r77, %r15, 31; add.s32 %r78, %r15, %r77; shr.s32 %r79, %r78, 1; cvt.s64.s32 %rd48, %r79; add.s64 %rd46, %rd35, %rd48; // begin inline asm createpolicy.fractional.L2::evict_first.L2::evict_first.b64 %rd45, 1.0; // end inline asm // begin inline asm ld.global.L1::evict_last.L2::cache_hint.u8 %rs83, [%rd46], %rd45; // end inline asm cvt.u32.u16 %r80, %rs83; and.b32 %r81, %r80, 255; shl.b32 %r82, %r14, 2; and.b32 %r83, %r82, 4; shr.u32 %r84, %r81, %r83; cvt.u16.u32 %rs84, %r84; and.b16 %rs787, %rs84, 15; $L__BB0_4: shl.b32 %r16, %r761, 5; setp.ge.s32 %p3, %r16, %r61; @%p3 bra $L__BB0_8; setp.eq.s16 %p4, %rs73, 0; shr.u16 %rs86, %rs787, 3; and.b16 %rs87, %rs86, 1; setp.eq.b16 %p5, %rs87, 1; and.pred %p6, %p4, %p5; selp.b16 %rs88, -16, 0, %p6; or.b16 %rs89, %rs88, %rs787; cvt.s16.s8 %rs90, %rs89; cvt.rn.f32.s16 %f10, %rs90; mul.wide.s32 %rd49, %r16, 2; add.s64 %rd7, %rd3, %rd49; ld.global.v4.u32 {%r85, %r86, %r87, %r88}, [%rd7]; mul.wide.s32 %rd50, %r61, 2; add.s64 %rd51, %rd7, %rd50; ld.global.v4.u32 {%r89, %r90, %r91, %r92}, [%rd51]; add.s32 %r93, %r16, %r61; add.s32 %r94, %r93, %r61; mul.wide.s32 %rd52, %r94, 2; add.s64 %rd8, %rd3, %rd52; ld.global.v4.u32 {%r95, %r96, %r97, %r98}, [%rd8]; add.s64 %rd53, %rd8, %rd50; ld.global.v4.u32 {%r99, %r100, %r101, %r102}, [%rd53]; add.s64 %rd54, %rd53, %rd50; ld.global.v4.u32 {%r103, %r104, %r105, %r106}, [%rd54]; add.s64 %rd55, %rd54, %rd50; ld.global.v4.u32 {%r107, %r108, %r109, %r110}, [%rd55]; add.s64 %rd56, %rd55, %rd50; ld.global.v4.u32 {%r111, %r112, %r113, %r114}, [%rd56]; add.s64 %rd57, %rd56, %rd50; ld.global.v4.u32 {%r115, %r116, %r117, %r118}, [%rd57]; cvt.u16.u32 %rs5, %r68; and.b16 %rs6, %rs5, 15; shr.u32 %r119, %r68, 4; cvt.u16.u32 %rs7, %r119; and.b16 %rs8, %rs7, 15; shr.u32 %r120, %r68, 8; cvt.u16.u32 %rs9, %r120; and.b16 %rs10, %rs9, 15; shr.u32 %r121, %r68, 12; cvt.u16.u32 %rs11, %r121; and.b16 %rs12, %rs11, 15; shr.u32 %r122, %r68, 16; cvt.u16.u32 %rs13, %r122; and.b16 %rs14, %rs13, 15; shr.u32 %r123, %r68, 20; cvt.u16.u32 %rs15, %r123; and.b16 %rs16, %rs15, 15; shr.u32 %r124, %r68, 24; cvt.u16.u32 %rs17, %r124; and.b16 %rs18, %rs17, 15; shr.u32 %r125, %r68, 28; cvt.u16.u32 %rs19, %r125; add.s32 %r126, %r93, 8; mul.wide.s32 %rd58, %r126, 2; add.s64 %rd9, %rd3, %rd58; add.s32 %r127, %r126, %r61; add.s32 %r128, %r127, %r61; mul.wide.s32 %rd59, %r128, 2; add.s64 %rd10, %rd3, %rd59; add.s32 %r129, %r128, %r61; mul.wide.s32 %rd60, %r129, 2; add.s64 %rd11, %rd3, %rd60; add.s32 %r130, %r129, %r61; mul.wide.s32 %rd61, %r130, 2; add.s64 %rd12, %rd3, %rd61; add.s32 %r131, %r130, %r61; mul.wide.s32 %rd62, %r131, 2; add.s64 %rd13, %rd3, %rd62; add.s32 %r132, %r131, %r61; mul.wide.s32 %rd63, %r132, 2; add.s64 %rd14, %rd3, %rd63; cvt.u16.u32 %rs20, %r69; and.b16 %rs21, %rs20, 15; shr.u32 %r133, %r69, 4; cvt.u16.u32 %rs22, %r133; and.b16 %rs23, %rs22, 15; shr.u32 %r134, %r69, 8; cvt.u16.u32 %rs24, %r134; and.b16 %rs25, %rs24, 15; shr.u32 %r135, %r69, 12; cvt.u16.u32 %rs26, %r135; and.b16 %rs27, %rs26, 15; shr.u32 %r136, %r69, 16; cvt.u16.u32 %rs28, %r136; and.b16 %rs29, %rs28, 15; shr.u32 %r137, %r69, 20; cvt.u16.u32 %rs30, %r137; and.b16 %rs31, %rs30, 15; shr.u32 %r138, %r69, 24; cvt.u16.u32 %rs32, %r138; and.b16 %rs33, %rs32, 15; shr.u32 %r139, %r69, 28; cvt.u16.u32 %rs34, %r139; cvt.u16.u32 %rs35, %r70; and.b16 %rs36, %rs35, 15; shr.u32 %r140, %r70, 4; cvt.u16.u32 %rs37, %r140; and.b16 %rs38, %rs37, 15; shr.u32 %r141, %r70, 8; cvt.u16.u32 %rs39, %r141; and.b16 %rs40, %rs39, 15; shr.u32 %r142, %r70, 12; cvt.u16.u32 %rs41, %r142; and.b16 %rs42, %rs41, 15; shr.u32 %r143, %r70, 16; cvt.u16.u32 %rs43, %r143; and.b16 %rs44, %rs43, 15; shr.u32 %r144, %r70, 20; cvt.u16.u32 %rs45, %r144; and.b16 %rs46, %rs45, 15; shr.u32 %r145, %r70, 24; cvt.u16.u32 %rs47, %r145; and.b16 %rs48, %rs47, 15; shr.u32 %r146, %r70, 28; cvt.u16.u32 %rs49, %r146; cvt.u16.u32 %rs50, %r71; and.b16 %rs51, %rs50, 15; shr.u32 %r147, %r71, 4; cvt.u16.u32 %rs52, %r147; and.b16 %rs53, %rs52, 15; shr.u32 %r148, %r71, 8; cvt.u16.u32 %rs54, %r148; and.b16 %rs55, %rs54, 15; shr.u32 %r149, %r71, 12; cvt.u16.u32 %rs56, %r149; and.b16 %rs57, %rs56, 15; shr.u32 %r150, %r71, 16; cvt.u16.u32 %rs58, %r150; and.b16 %rs59, %rs58, 15; shr.u32 %r151, %r71, 20; cvt.u16.u32 %rs60, %r151; and.b16 %rs61, %rs60, 15; shr.u32 %r152, %r71, 24; cvt.u16.u32 %rs62, %r152; and.b16 %rs63, %rs62, 15; shr.u32 %r153, %r71, 28; cvt.u16.u32 %rs64, %r153; @%p4 bra $L__BB0_7; cvt.rn.f32.s16 %f362, %rs6; sub.ftz.f32 %f363, %f362, %f10; mul.ftz.f32 %f364, %f105, %f363; mov.b32 {%rs91, %rs99}, %r85; // begin inline asm { cvt.f32.f16 %f106, %rs91;} // end inline asm fma.rn.ftz.f32 %f365, %f364, %f106, %f1497; mov.b32 {%rs92, %rs100}, %r89; // begin inline asm { cvt.f32.f16 %f107, %rs92;} // end inline asm fma.rn.ftz.f32 %f366, %f364, %f107, %f1496; mov.b32 {%rs93, %rs101}, %r95; // begin inline asm { cvt.f32.f16 %f108, %rs93;} // end inline asm fma.rn.ftz.f32 %f367, %f364, %f108, %f1495; mov.b32 {%rs94, %rs102}, %r99; // begin inline asm { cvt.f32.f16 %f109, %rs94;} // end inline asm fma.rn.ftz.f32 %f368, %f364, %f109, %f1494; mov.b32 {%rs95, %rs103}, %r103; // begin inline asm { cvt.f32.f16 %f110, %rs95;} // end inline asm fma.rn.ftz.f32 %f369, %f364, %f110, %f1493; mov.b32 {%rs96, %rs104}, %r107; // begin inline asm { cvt.f32.f16 %f111, %rs96;} // end inline asm fma.rn.ftz.f32 %f370, %f364, %f111, %f1492; mov.b32 {%rs97, %rs105}, %r111; // begin inline asm { cvt.f32.f16 %f112, %rs97;} // end inline asm fma.rn.ftz.f32 %f371, %f364, %f112, %f1491; mov.b32 {%rs98, %rs106}, %r115; // begin inline asm { cvt.f32.f16 %f113, %rs98;} // end inline asm fma.rn.ftz.f32 %f372, %f364, %f113, %f1490; cvt.rn.f32.s16 %f373, %rs8; sub.ftz.f32 %f374, %f373, %f10; mul.ftz.f32 %f375, %f105, %f374; // begin inline asm { cvt.f32.f16 %f114, %rs99;} // end inline asm fma.rn.ftz.f32 %f376, %f375, %f114, %f365; // begin inline asm { cvt.f32.f16 %f115, %rs100;} // end inline asm fma.rn.ftz.f32 %f377, %f375, %f115, %f366; // begin inline asm { cvt.f32.f16 %f116, %rs101;} // end inline asm fma.rn.ftz.f32 %f378, %f375, %f116, %f367; // begin inline asm { cvt.f32.f16 %f117, %rs102;} // end inline asm fma.rn.ftz.f32 %f379, %f375, %f117, %f368; // begin inline asm { cvt.f32.f16 %f118, %rs103;} // end inline asm fma.rn.ftz.f32 %f380, %f375, %f118, %f369; // begin inline asm { cvt.f32.f16 %f119, %rs104;} // end inline asm fma.rn.ftz.f32 %f381, %f375, %f119, %f370; // begin inline asm { cvt.f32.f16 %f120, %rs105;} // end inline asm fma.rn.ftz.f32 %f382, %f375, %f120, %f371; // begin inline asm { cvt.f32.f16 %f121, %rs106;} // end inline asm fma.rn.ftz.f32 %f383, %f375, %f121, %f372; cvt.rn.f32.s16 %f384, %rs10; sub.ftz.f32 %f385, %f384, %f10; mul.ftz.f32 %f386, %f105, %f385; mov.b32 {%rs107, %rs115}, %r86; // begin inline asm { cvt.f32.f16 %f122, %rs107;} // end inline asm fma.rn.ftz.f32 %f387, %f386, %f122, %f376; mov.b32 {%rs108, %rs116}, %r90; // begin inline asm { cvt.f32.f16 %f123, %rs108;} // end inline asm fma.rn.ftz.f32 %f388, %f386, %f123, %f377; mov.b32 {%rs109, %rs117}, %r96; // begin inline asm { cvt.f32.f16 %f124, %rs109;} // end inline asm fma.rn.ftz.f32 %f389, %f386, %f124, %f378; mov.b32 {%rs110, %rs118}, %r100; // begin inline asm { cvt.f32.f16 %f125, %rs110;} // end inline asm fma.rn.ftz.f32 %f390, %f386, %f125, %f379; mov.b32 {%rs111, %rs119}, %r104; // begin inline asm { cvt.f32.f16 %f126, %rs111;} // end inline asm fma.rn.ftz.f32 %f391, %f386, %f126, %f380; mov.b32 {%rs112, %rs120}, %r108; // begin inline asm { cvt.f32.f16 %f127, %rs112;} // end inline asm fma.rn.ftz.f32 %f392, %f386, %f127, %f381; mov.b32 {%rs113, %rs121}, %r112; // begin inline asm { cvt.f32.f16 %f128, %rs113;} // end inline asm fma.rn.ftz.f32 %f393, %f386, %f128, %f382; mov.b32 {%rs114, %rs122}, %r116; // begin inline asm { cvt.f32.f16 %f129, %rs114;} // end inline asm fma.rn.ftz.f32 %f394, %f386, %f129, %f383; cvt.rn.f32.s16 %f395, %rs12; sub.ftz.f32 %f396, %f395, %f10; mul.ftz.f32 %f397, %f105, %f396; // begin inline asm { cvt.f32.f16 %f130, %rs115;} // end inline asm fma.rn.ftz.f32 %f398, %f397, %f130, %f387; // begin inline asm { cvt.f32.f16 %f131, %rs116;} // end inline asm fma.rn.ftz.f32 %f399, %f397, %f131, %f388; // begin inline asm { cvt.f32.f16 %f132, %rs117;} // end inline asm fma.rn.ftz.f32 %f400, %f397, %f132, %f389; // begin inline asm { cvt.f32.f16 %f133, %rs118;} // end inline asm fma.rn.ftz.f32 %f401, %f397, %f133, %f390; // begin inline asm { cvt.f32.f16 %f134, %rs119;} // end inline asm fma.rn.ftz.f32 %f402, %f397, %f134, %f391; // begin inline asm { cvt.f32.f16 %f135, %rs120;} // end inline asm fma.rn.ftz.f32 %f403, %f397, %f135, %f392; // begin inline asm { cvt.f32.f16 %f136, %rs121;} // end inline asm fma.rn.ftz.f32 %f404, %f397, %f136, %f393; // begin inline asm { cvt.f32.f16 %f137, %rs122;} // end inline asm fma.rn.ftz.f32 %f405, %f397, %f137, %f394; cvt.rn.f32.s16 %f406, %rs14; sub.ftz.f32 %f407, %f406, %f10; mul.ftz.f32 %f408, %f105, %f407; mov.b32 {%rs123, %rs131}, %r87; // begin inline asm { cvt.f32.f16 %f138, %rs123;} // end inline asm fma.rn.ftz.f32 %f409, %f408, %f138, %f398; mov.b32 {%rs124, %rs132}, %r91; // begin inline asm { cvt.f32.f16 %f139, %rs124;} // end inline asm fma.rn.ftz.f32 %f410, %f408, %f139, %f399; mov.b32 {%rs125, %rs133}, %r97; // begin inline asm { cvt.f32.f16 %f140, %rs125;} // end inline asm fma.rn.ftz.f32 %f411, %f408, %f140, %f400; mov.b32 {%rs126, %rs134}, %r101; // begin inline asm { cvt.f32.f16 %f141, %rs126;} // end inline asm fma.rn.ftz.f32 %f412, %f408, %f141, %f401; mov.b32 {%rs127, %rs135}, %r105; // begin inline asm { cvt.f32.f16 %f142, %rs127;} // end inline asm fma.rn.ftz.f32 %f413, %f408, %f142, %f402; mov.b32 {%rs128, %rs136}, %r109; // begin inline asm { cvt.f32.f16 %f143, %rs128;} // end inline asm fma.rn.ftz.f32 %f414, %f408, %f143, %f403; mov.b32 {%rs129, %rs137}, %r113; // begin inline asm { cvt.f32.f16 %f144, %rs129;} // end inline asm fma.rn.ftz.f32 %f415, %f408, %f144, %f404; mov.b32 {%rs130, %rs138}, %r117; // begin inline asm { cvt.f32.f16 %f145, %rs130;} // end inline asm fma.rn.ftz.f32 %f416, %f408, %f145, %f405; cvt.rn.f32.s16 %f417, %rs16; sub.ftz.f32 %f418, %f417, %f10; mul.ftz.f32 %f419, %f105, %f418; // begin inline asm { cvt.f32.f16 %f146, %rs131;} // end inline asm fma.rn.ftz.f32 %f420, %f419, %f146, %f409; // begin inline asm { cvt.f32.f16 %f147, %rs132;} // end inline asm fma.rn.ftz.f32 %f421, %f419, %f147, %f410; // begin inline asm { cvt.f32.f16 %f148, %rs133;} // end inline asm fma.rn.ftz.f32 %f422, %f419, %f148, %f411; // begin inline asm { cvt.f32.f16 %f149, %rs134;} // end inline asm fma.rn.ftz.f32 %f423, %f419, %f149, %f412; // begin inline asm { cvt.f32.f16 %f150, %rs135;} // end inline asm fma.rn.ftz.f32 %f424, %f419, %f150, %f413; // begin inline asm { cvt.f32.f16 %f151, %rs136;} // end inline asm fma.rn.ftz.f32 %f425, %f419, %f151, %f414; // begin inline asm { cvt.f32.f16 %f152, %rs137;} // end inline asm fma.rn.ftz.f32 %f426, %f419, %f152, %f415; // begin inline asm { cvt.f32.f16 %f153, %rs138;} // end inline asm fma.rn.ftz.f32 %f427, %f419, %f153, %f416; cvt.rn.f32.s16 %f428, %rs18; sub.ftz.f32 %f429, %f428, %f10; mul.ftz.f32 %f430, %f105, %f429; mov.b32 {%rs139, %rs147}, %r88; // begin inline asm { cvt.f32.f16 %f154, %rs139;} // end inline asm fma.rn.ftz.f32 %f431, %f430, %f154, %f420; mov.b32 {%rs140, %rs148}, %r92; // begin inline asm { cvt.f32.f16 %f155, %rs140;} // end inline asm fma.rn.ftz.f32 %f432, %f430, %f155, %f421; mov.b32 {%rs141, %rs149}, %r98; // begin inline asm { cvt.f32.f16 %f156, %rs141;} // end inline asm fma.rn.ftz.f32 %f433, %f430, %f156, %f422; mov.b32 {%rs142, %rs150}, %r102; // begin inline asm { cvt.f32.f16 %f157, %rs142;} // end inline asm fma.rn.ftz.f32 %f434, %f430, %f157, %f423; mov.b32 {%rs143, %rs151}, %r106; // begin inline asm { cvt.f32.f16 %f158, %rs143;} // end inline asm fma.rn.ftz.f32 %f435, %f430, %f158, %f424; mov.b32 {%rs144, %rs152}, %r110; // begin inline asm { cvt.f32.f16 %f159, %rs144;} // end inline asm fma.rn.ftz.f32 %f436, %f430, %f159, %f425; mov.b32 {%rs145, %rs153}, %r114; // begin inline asm { cvt.f32.f16 %f160, %rs145;} // end inline asm fma.rn.ftz.f32 %f437, %f430, %f160, %f426; mov.b32 {%rs146, %rs154}, %r118; // begin inline asm { cvt.f32.f16 %f161, %rs146;} // end inline asm fma.rn.ftz.f32 %f438, %f430, %f161, %f427; cvt.rn.f32.s16 %f439, %rs19; sub.ftz.f32 %f440, %f439, %f10; mul.ftz.f32 %f441, %f105, %f440; // begin inline asm { cvt.f32.f16 %f162, %rs147;} // end inline asm fma.rn.ftz.f32 %f442, %f441, %f162, %f431; // begin inline asm { cvt.f32.f16 %f163, %rs148;} // end inline asm fma.rn.ftz.f32 %f443, %f441, %f163, %f432; // begin inline asm { cvt.f32.f16 %f164, %rs149;} // end inline asm fma.rn.ftz.f32 %f444, %f441, %f164, %f433; // begin inline asm { cvt.f32.f16 %f165, %rs150;} // end inline asm fma.rn.ftz.f32 %f445, %f441, %f165, %f434; // begin inline asm { cvt.f32.f16 %f166, %rs151;} // end inline asm fma.rn.ftz.f32 %f446, %f441, %f166, %f435; // begin inline asm { cvt.f32.f16 %f167, %rs152;} // end inline asm fma.rn.ftz.f32 %f447, %f441, %f167, %f436; // begin inline asm { cvt.f32.f16 %f168, %rs153;} // end inline asm fma.rn.ftz.f32 %f448, %f441, %f168, %f437; // begin inline asm { cvt.f32.f16 %f169, %rs154;} // end inline asm fma.rn.ftz.f32 %f449, %f441, %f169, %f438; ld.global.v4.u32 {%r154, %r155, %r156, %r157}, [%rd7+16]; ld.global.v4.u32 {%r162, %r163, %r164, %r165}, [%rd9]; ld.global.v4.u32 {%r170, %r171, %r172, %r173}, [%rd8+16]; ld.global.v4.u32 {%r178, %r179, %r180, %r181}, [%rd10]; ld.global.v4.u32 {%r186, %r187, %r188, %r189}, [%rd11]; ld.global.v4.u32 {%r194, %r195, %r196, %r197}, [%rd12]; ld.global.v4.u32 {%r202, %r203, %r204, %r205}, [%rd13]; ld.global.v4.u32 {%r210, %r211, %r212, %r213}, [%rd14]; cvt.rn.f32.s16 %f450, %rs21; sub.ftz.f32 %f451, %f450, %f10; mul.ftz.f32 %f452, %f105, %f451; mov.b32 {%rs155, %rs163}, %r154; // begin inline asm { cvt.f32.f16 %f170, %rs155;} // end inline asm fma.rn.ftz.f32 %f453, %f452, %f170, %f442; mov.b32 {%rs156, %rs164}, %r162; // begin inline asm { cvt.f32.f16 %f171, %rs156;} // end inline asm fma.rn.ftz.f32 %f454, %f452, %f171, %f443; mov.b32 {%rs157, %rs165}, %r170; // begin inline asm { cvt.f32.f16 %f172, %rs157;} // end inline asm fma.rn.ftz.f32 %f455, %f452, %f172, %f444; mov.b32 {%rs158, %rs166}, %r178; // begin inline asm { cvt.f32.f16 %f173, %rs158;} // end inline asm fma.rn.ftz.f32 %f456, %f452, %f173, %f445; mov.b32 {%rs159, %rs167}, %r186; // begin inline asm { cvt.f32.f16 %f174, %rs159;} // end inline asm fma.rn.ftz.f32 %f457, %f452, %f174, %f446; mov.b32 {%rs160, %rs168}, %r194; // begin inline asm { cvt.f32.f16 %f175, %rs160;} // end inline asm fma.rn.ftz.f32 %f458, %f452, %f175, %f447; mov.b32 {%rs161, %rs169}, %r202; // begin inline asm { cvt.f32.f16 %f176, %rs161;} // end inline asm fma.rn.ftz.f32 %f459, %f452, %f176, %f448; mov.b32 {%rs162, %rs170}, %r210; // begin inline asm { cvt.f32.f16 %f177, %rs162;} // end inline asm fma.rn.ftz.f32 %f460, %f452, %f177, %f449; cvt.rn.f32.s16 %f461, %rs23; sub.ftz.f32 %f462, %f461, %f10; mul.ftz.f32 %f463, %f105, %f462; // begin inline asm { cvt.f32.f16 %f178, %rs163;} // end inline asm fma.rn.ftz.f32 %f464, %f463, %f178, %f453; // begin inline asm { cvt.f32.f16 %f179, %rs164;} // end inline asm fma.rn.ftz.f32 %f465, %f463, %f179, %f454; // begin inline asm { cvt.f32.f16 %f180, %rs165;} // end inline asm fma.rn.ftz.f32 %f466, %f463, %f180, %f455; // begin inline asm { cvt.f32.f16 %f181, %rs166;} // end inline asm fma.rn.ftz.f32 %f467, %f463, %f181, %f456; // begin inline asm { cvt.f32.f16 %f182, %rs167;} // end inline asm fma.rn.ftz.f32 %f468, %f463, %f182, %f457; // begin inline asm { cvt.f32.f16 %f183, %rs168;} // end inline asm fma.rn.ftz.f32 %f469, %f463, %f183, %f458; // begin inline asm { cvt.f32.f16 %f184, %rs169;} // end inline asm fma.rn.ftz.f32 %f470, %f463, %f184, %f459; // begin inline asm { cvt.f32.f16 %f185, %rs170;} // end inline asm fma.rn.ftz.f32 %f471, %f463, %f185, %f460; cvt.rn.f32.s16 %f472, %rs25; sub.ftz.f32 %f473, %f472, %f10; mul.ftz.f32 %f474, %f105, %f473; mov.b32 {%rs171, %rs179}, %r155; // begin inline asm { cvt.f32.f16 %f186, %rs171;} // end inline asm fma.rn.ftz.f32 %f475, %f474, %f186, %f464; mov.b32 {%rs172, %rs180}, %r163; // begin inline asm { cvt.f32.f16 %f187, %rs172;} // end inline asm fma.rn.ftz.f32 %f476, %f474, %f187, %f465; mov.b32 {%rs173, %rs181}, %r171; // begin inline asm { cvt.f32.f16 %f188, %rs173;} // end inline asm fma.rn.ftz.f32 %f477, %f474, %f188, %f466; mov.b32 {%rs174, %rs182}, %r179; // begin inline asm { cvt.f32.f16 %f189, %rs174;} // end inline asm fma.rn.ftz.f32 %f478, %f474, %f189, %f467; mov.b32 {%rs175, %rs183}, %r187; // begin inline asm { cvt.f32.f16 %f190, %rs175;} // end inline asm fma.rn.ftz.f32 %f479, %f474, %f190, %f468; mov.b32 {%rs176, %rs184}, %r195; // begin inline asm { cvt.f32.f16 %f191, %rs176;} // end inline asm fma.rn.ftz.f32 %f480, %f474, %f191, %f469; mov.b32 {%rs177, %rs185}, %r203; // begin inline asm { cvt.f32.f16 %f192, %rs177;} // end inline asm fma.rn.ftz.f32 %f481, %f474, %f192, %f470; mov.b32 {%rs178, %rs186}, %r211; // begin inline asm { cvt.f32.f16 %f193, %rs178;} // end inline asm fma.rn.ftz.f32 %f482, %f474, %f193, %f471; cvt.rn.f32.s16 %f483, %rs27; sub.ftz.f32 %f484, %f483, %f10; mul.ftz.f32 %f485, %f105, %f484; // begin inline asm { cvt.f32.f16 %f194, %rs179;} // end inline asm fma.rn.ftz.f32 %f486, %f485, %f194, %f475; // begin inline asm { cvt.f32.f16 %f195, %rs180;} // end inline asm fma.rn.ftz.f32 %f487, %f485, %f195, %f476; // begin inline asm { cvt.f32.f16 %f196, %rs181;} // end inline asm fma.rn.ftz.f32 %f488, %f485, %f196, %f477; // begin inline asm { cvt.f32.f16 %f197, %rs182;} // end inline asm fma.rn.ftz.f32 %f489, %f485, %f197, %f478; // begin inline asm { cvt.f32.f16 %f198, %rs183;} // end inline asm fma.rn.ftz.f32 %f490, %f485, %f198, %f479; // begin inline asm { cvt.f32.f16 %f199, %rs184;} // end inline asm fma.rn.ftz.f32 %f491, %f485, %f199, %f480; // begin inline asm { cvt.f32.f16 %f200, %rs185;} // end inline asm fma.rn.ftz.f32 %f492, %f485, %f200, %f481; // begin inline asm { cvt.f32.f16 %f201, %rs186;} // end inline asm fma.rn.ftz.f32 %f493, %f485, %f201, %f482; cvt.rn.f32.s16 %f494, %rs29; sub.ftz.f32 %f495, %f494, %f10; mul.ftz.f32 %f496, %f105, %f495; mov.b32 {%rs187, %rs195}, %r156; // begin inline asm { cvt.f32.f16 %f202, %rs187;} // end inline asm fma.rn.ftz.f32 %f497, %f496, %f202, %f486; mov.b32 {%rs188, %rs196}, %r164; // begin inline asm { cvt.f32.f16 %f203, %rs188;} // end inline asm fma.rn.ftz.f32 %f498, %f496, %f203, %f487; mov.b32 {%rs189, %rs197}, %r172; // begin inline asm { cvt.f32.f16 %f204, %rs189;} // end inline asm fma.rn.ftz.f32 %f499, %f496, %f204, %f488; mov.b32 {%rs190, %rs198}, %r180; // begin inline asm { cvt.f32.f16 %f205, %rs190;} // end inline asm fma.rn.ftz.f32 %f500, %f496, %f205, %f489; mov.b32 {%rs191, %rs199}, %r188; // begin inline asm { cvt.f32.f16 %f206, %rs191;} // end inline asm fma.rn.ftz.f32 %f501, %f496, %f206, %f490; mov.b32 {%rs192, %rs200}, %r196; // begin inline asm { cvt.f32.f16 %f207, %rs192;} // end inline asm fma.rn.ftz.f32 %f502, %f496, %f207, %f491; mov.b32 {%rs193, %rs201}, %r204; // begin inline asm { cvt.f32.f16 %f208, %rs193;} // end inline asm fma.rn.ftz.f32 %f503, %f496, %f208, %f492; mov.b32 {%rs194, %rs202}, %r212; // begin inline asm { cvt.f32.f16 %f209, %rs194;} // end inline asm fma.rn.ftz.f32 %f504, %f496, %f209, %f493; cvt.rn.f32.s16 %f505, %rs31; sub.ftz.f32 %f506, %f505, %f10; mul.ftz.f32 %f507, %f105, %f506; // begin inline asm { cvt.f32.f16 %f210, %rs195;} // end inline asm fma.rn.ftz.f32 %f508, %f507, %f210, %f497; // begin inline asm { cvt.f32.f16 %f211, %rs196;} // end inline asm fma.rn.ftz.f32 %f509, %f507, %f211, %f498; // begin inline asm { cvt.f32.f16 %f212, %rs197;} // end inline asm fma.rn.ftz.f32 %f510, %f507, %f212, %f499; // begin inline asm { cvt.f32.f16 %f213, %rs198;} // end inline asm fma.rn.ftz.f32 %f511, %f507, %f213, %f500; // begin inline asm { cvt.f32.f16 %f214, %rs199;} // end inline asm fma.rn.ftz.f32 %f512, %f507, %f214, %f501; // begin inline asm { cvt.f32.f16 %f215, %rs200;} // end inline asm fma.rn.ftz.f32 %f513, %f507, %f215, %f502; // begin inline asm { cvt.f32.f16 %f216, %rs201;} // end inline asm fma.rn.ftz.f32 %f514, %f507, %f216, %f503; // begin inline asm { cvt.f32.f16 %f217, %rs202;} // end inline asm fma.rn.ftz.f32 %f515, %f507, %f217, %f504; cvt.rn.f32.s16 %f516, %rs33; sub.ftz.f32 %f517, %f516, %f10; mul.ftz.f32 %f518, %f105, %f517; mov.b32 {%rs203, %rs211}, %r157; // begin inline asm { cvt.f32.f16 %f218, %rs203;} // end inline asm fma.rn.ftz.f32 %f519, %f518, %f218, %f508; mov.b32 {%rs204, %rs212}, %r165; // begin inline asm { cvt.f32.f16 %f219, %rs204;} // end inline asm fma.rn.ftz.f32 %f520, %f518, %f219, %f509; mov.b32 {%rs205, %rs213}, %r173; // begin inline asm { cvt.f32.f16 %f220, %rs205;} // end inline asm fma.rn.ftz.f32 %f521, %f518, %f220, %f510; mov.b32 {%rs206, %rs214}, %r181; // begin inline asm { cvt.f32.f16 %f221, %rs206;} // end inline asm fma.rn.ftz.f32 %f522, %f518, %f221, %f511; mov.b32 {%rs207, %rs215}, %r189; // begin inline asm { cvt.f32.f16 %f222, %rs207;} // end inline asm fma.rn.ftz.f32 %f523, %f518, %f222, %f512; mov.b32 {%rs208, %rs216}, %r197; // begin inline asm { cvt.f32.f16 %f223, %rs208;} // end inline asm fma.rn.ftz.f32 %f524, %f518, %f223, %f513; mov.b32 {%rs209, %rs217}, %r205; // begin inline asm { cvt.f32.f16 %f224, %rs209;} // end inline asm fma.rn.ftz.f32 %f525, %f518, %f224, %f514; mov.b32 {%rs210, %rs218}, %r213; // begin inline asm { cvt.f32.f16 %f225, %rs210;} // end inline asm fma.rn.ftz.f32 %f526, %f518, %f225, %f515; cvt.rn.f32.s16 %f527, %rs34; sub.ftz.f32 %f528, %f527, %f10; mul.ftz.f32 %f529, %f105, %f528; // begin inline asm { cvt.f32.f16 %f226, %rs211;} // end inline asm fma.rn.ftz.f32 %f530, %f529, %f226, %f519; // begin inline asm { cvt.f32.f16 %f227, %rs212;} // end inline asm fma.rn.ftz.f32 %f531, %f529, %f227, %f520; // begin inline asm { cvt.f32.f16 %f228, %rs213;} // end inline asm fma.rn.ftz.f32 %f532, %f529, %f228, %f521; // begin inline asm { cvt.f32.f16 %f229, %rs214;} // end inline asm fma.rn.ftz.f32 %f533, %f529, %f229, %f522; // begin inline asm { cvt.f32.f16 %f230, %rs215;} // end inline asm fma.rn.ftz.f32 %f534, %f529, %f230, %f523; // begin inline asm { cvt.f32.f16 %f231, %rs216;} // end inline asm fma.rn.ftz.f32 %f535, %f529, %f231, %f524; // begin inline asm { cvt.f32.f16 %f232, %rs217;} // end inline asm fma.rn.ftz.f32 %f536, %f529, %f232, %f525; // begin inline asm { cvt.f32.f16 %f233, %rs218;} // end inline asm fma.rn.ftz.f32 %f537, %f529, %f233, %f526; ld.global.v4.u32 {%r218, %r219, %r220, %r221}, [%rd7+32]; ld.global.v4.u32 {%r226, %r227, %r228, %r229}, [%rd9+16]; ld.global.v4.u32 {%r234, %r235, %r236, %r237}, [%rd8+32]; ld.global.v4.u32 {%r242, %r243, %r244, %r245}, [%rd10+16]; ld.global.v4.u32 {%r250, %r251, %r252, %r253}, [%rd11+16]; ld.global.v4.u32 {%r258, %r259, %r260, %r261}, [%rd12+16]; ld.global.v4.u32 {%r266, %r267, %r268, %r269}, [%rd13+16]; ld.global.v4.u32 {%r274, %r275, %r276, %r277}, [%rd14+16]; cvt.rn.f32.s16 %f538, %rs36; sub.ftz.f32 %f539, %f538, %f10; mul.ftz.f32 %f540, %f105, %f539; mov.b32 {%rs219, %rs227}, %r218; // begin inline asm { cvt.f32.f16 %f234, %rs219;} // end inline asm fma.rn.ftz.f32 %f541, %f540, %f234, %f530; mov.b32 {%rs220, %rs228}, %r226; // begin inline asm { cvt.f32.f16 %f235, %rs220;} // end inline asm fma.rn.ftz.f32 %f542, %f540, %f235, %f531; mov.b32 {%rs221, %rs229}, %r234; // begin inline asm { cvt.f32.f16 %f236, %rs221;} // end inline asm fma.rn.ftz.f32 %f543, %f540, %f236, %f532; mov.b32 {%rs222, %rs230}, %r242; // begin inline asm { cvt.f32.f16 %f237, %rs222;} // end inline asm fma.rn.ftz.f32 %f544, %f540, %f237, %f533; mov.b32 {%rs223, %rs231}, %r250; // begin inline asm { cvt.f32.f16 %f238, %rs223;} // end inline asm fma.rn.ftz.f32 %f545, %f540, %f238, %f534; mov.b32 {%rs224, %rs232}, %r258; // begin inline asm { cvt.f32.f16 %f239, %rs224;} // end inline asm fma.rn.ftz.f32 %f546, %f540, %f239, %f535; mov.b32 {%rs225, %rs233}, %r266; // begin inline asm { cvt.f32.f16 %f240, %rs225;} // end inline asm fma.rn.ftz.f32 %f547, %f540, %f240, %f536; mov.b32 {%rs226, %rs234}, %r274; // begin inline asm { cvt.f32.f16 %f241, %rs226;} // end inline asm fma.rn.ftz.f32 %f548, %f540, %f241, %f537; cvt.rn.f32.s16 %f549, %rs38; sub.ftz.f32 %f550, %f549, %f10; mul.ftz.f32 %f551, %f105, %f550; // begin inline asm { cvt.f32.f16 %f242, %rs227;} // end inline asm fma.rn.ftz.f32 %f552, %f551, %f242, %f541; // begin inline asm { cvt.f32.f16 %f243, %rs228;} // end inline asm fma.rn.ftz.f32 %f553, %f551, %f243, %f542; // begin inline asm { cvt.f32.f16 %f244, %rs229;} // end inline asm fma.rn.ftz.f32 %f554, %f551, %f244, %f543; // begin inline asm { cvt.f32.f16 %f245, %rs230;} // end inline asm fma.rn.ftz.f32 %f555, %f551, %f245, %f544; // begin inline asm { cvt.f32.f16 %f246, %rs231;} // end inline asm fma.rn.ftz.f32 %f556, %f551, %f246, %f545; // begin inline asm { cvt.f32.f16 %f247, %rs232;} // end inline asm fma.rn.ftz.f32 %f557, %f551, %f247, %f546; // begin inline asm { cvt.f32.f16 %f248, %rs233;} // end inline asm fma.rn.ftz.f32 %f558, %f551, %f248, %f547; // begin inline asm { cvt.f32.f16 %f249, %rs234;} // end inline asm fma.rn.ftz.f32 %f559, %f551, %f249, %f548; cvt.rn.f32.s16 %f560, %rs40; sub.ftz.f32 %f561, %f560, %f10; mul.ftz.f32 %f562, %f105, %f561; mov.b32 {%rs235, %rs243}, %r219; // begin inline asm { cvt.f32.f16 %f250, %rs235;} // end inline asm fma.rn.ftz.f32 %f563, %f562, %f250, %f552; mov.b32 {%rs236, %rs244}, %r227; // begin inline asm { cvt.f32.f16 %f251, %rs236;} // end inline asm fma.rn.ftz.f32 %f564, %f562, %f251, %f553; mov.b32 {%rs237, %rs245}, %r235; // begin inline asm { cvt.f32.f16 %f252, %rs237;} // end inline asm fma.rn.ftz.f32 %f565, %f562, %f252, %f554; mov.b32 {%rs238, %rs246}, %r243; // begin inline asm { cvt.f32.f16 %f253, %rs238;} // end inline asm fma.rn.ftz.f32 %f566, %f562, %f253, %f555; mov.b32 {%rs239, %rs247}, %r251; // begin inline asm { cvt.f32.f16 %f254, %rs239;} // end inline asm fma.rn.ftz.f32 %f567, %f562, %f254, %f556; mov.b32 {%rs240, %rs248}, %r259; // begin inline asm { cvt.f32.f16 %f255, %rs240;} // end inline asm fma.rn.ftz.f32 %f568, %f562, %f255, %f557; mov.b32 {%rs241, %rs249}, %r267; // begin inline asm { cvt.f32.f16 %f256, %rs241;} // end inline asm fma.rn.ftz.f32 %f569, %f562, %f256, %f558; mov.b32 {%rs242, %rs250}, %r275; // begin inline asm { cvt.f32.f16 %f257, %rs242;} // end inline asm fma.rn.ftz.f32 %f570, %f562, %f257, %f559; cvt.rn.f32.s16 %f571, %rs42; sub.ftz.f32 %f572, %f571, %f10; mul.ftz.f32 %f573, %f105, %f572; // begin inline asm { cvt.f32.f16 %f258, %rs243;} // end inline asm fma.rn.ftz.f32 %f574, %f573, %f258, %f563; // begin inline asm { cvt.f32.f16 %f259, %rs244;} // end inline asm fma.rn.ftz.f32 %f575, %f573, %f259, %f564; // begin inline asm { cvt.f32.f16 %f260, %rs245;} // end inline asm fma.rn.ftz.f32 %f576, %f573, %f260, %f565; // begin inline asm { cvt.f32.f16 %f261, %rs246;} // end inline asm fma.rn.ftz.f32 %f577, %f573, %f261, %f566; // begin inline asm { cvt.f32.f16 %f262, %rs247;} // end inline asm fma.rn.ftz.f32 %f578, %f573, %f262, %f567; // begin inline asm { cvt.f32.f16 %f263, %rs248;} // end inline asm fma.rn.ftz.f32 %f579, %f573, %f263, %f568; // begin inline asm { cvt.f32.f16 %f264, %rs249;} // end inline asm fma.rn.ftz.f32 %f580, %f573, %f264, %f569; // begin inline asm { cvt.f32.f16 %f265, %rs250;} // end inline asm fma.rn.ftz.f32 %f581, %f573, %f265, %f570; cvt.rn.f32.s16 %f582, %rs44; sub.ftz.f32 %f583, %f582, %f10; mul.ftz.f32 %f584, %f105, %f583; mov.b32 {%rs251, %rs259}, %r220; // begin inline asm { cvt.f32.f16 %f266, %rs251;} // end inline asm fma.rn.ftz.f32 %f585, %f584, %f266, %f574; mov.b32 {%rs252, %rs260}, %r228; // begin inline asm { cvt.f32.f16 %f267, %rs252;} // end inline asm fma.rn.ftz.f32 %f586, %f584, %f267, %f575; mov.b32 {%rs253, %rs261}, %r236; // begin inline asm { cvt.f32.f16 %f268, %rs253;} // end inline asm fma.rn.ftz.f32 %f587, %f584, %f268, %f576; mov.b32 {%rs254, %rs262}, %r244; // begin inline asm { cvt.f32.f16 %f269, %rs254;} // end inline asm fma.rn.ftz.f32 %f588, %f584, %f269, %f577; mov.b32 {%rs255, %rs263}, %r252; // begin inline asm { cvt.f32.f16 %f270, %rs255;} // end inline asm fma.rn.ftz.f32 %f589, %f584, %f270, %f578; mov.b32 {%rs256, %rs264}, %r260; // begin inline asm { cvt.f32.f16 %f271, %rs256;} // end inline asm fma.rn.ftz.f32 %f590, %f584, %f271, %f579; mov.b32 {%rs257, %rs265}, %r268; // begin inline asm { cvt.f32.f16 %f272, %rs257;} // end inline asm fma.rn.ftz.f32 %f591, %f584, %f272, %f580; mov.b32 {%rs258, %rs266}, %r276; // begin inline asm { cvt.f32.f16 %f273, %rs258;} // end inline asm fma.rn.ftz.f32 %f592, %f584, %f273, %f581; cvt.rn.f32.s16 %f593, %rs46; sub.ftz.f32 %f594, %f593, %f10; mul.ftz.f32 %f595, %f105, %f594; // begin inline asm { cvt.f32.f16 %f274, %rs259;} // end inline asm fma.rn.ftz.f32 %f596, %f595, %f274, %f585; // begin inline asm { cvt.f32.f16 %f275, %rs260;} // end inline asm fma.rn.ftz.f32 %f597, %f595, %f275, %f586; // begin inline asm { cvt.f32.f16 %f276, %rs261;} // end inline asm fma.rn.ftz.f32 %f598, %f595, %f276, %f587; // begin inline asm { cvt.f32.f16 %f277, %rs262;} // end inline asm fma.rn.ftz.f32 %f599, %f595, %f277, %f588; // begin inline asm { cvt.f32.f16 %f278, %rs263;} // end inline asm fma.rn.ftz.f32 %f600, %f595, %f278, %f589; // begin inline asm { cvt.f32.f16 %f279, %rs264;} // end inline asm fma.rn.ftz.f32 %f601, %f595, %f279, %f590; // begin inline asm { cvt.f32.f16 %f280, %rs265;} // end inline asm fma.rn.ftz.f32 %f602, %f595, %f280, %f591; // begin inline asm { cvt.f32.f16 %f281, %rs266;} // end inline asm fma.rn.ftz.f32 %f603, %f595, %f281, %f592; cvt.rn.f32.s16 %f604, %rs48; sub.ftz.f32 %f605, %f604, %f10; mul.ftz.f32 %f606, %f105, %f605; mov.b32 {%rs267, %rs275}, %r221; // begin inline asm { cvt.f32.f16 %f282, %rs267;} // end inline asm fma.rn.ftz.f32 %f607, %f606, %f282, %f596; mov.b32 {%rs268, %rs276}, %r229; // begin inline asm { cvt.f32.f16 %f283, %rs268;} // end inline asm fma.rn.ftz.f32 %f608, %f606, %f283, %f597; mov.b32 {%rs269, %rs277}, %r237; // begin inline asm { cvt.f32.f16 %f284, %rs269;} // end inline asm fma.rn.ftz.f32 %f609, %f606, %f284, %f598; mov.b32 {%rs270, %rs278}, %r245; // begin inline asm { cvt.f32.f16 %f285, %rs270;} // end inline asm fma.rn.ftz.f32 %f610, %f606, %f285, %f599; mov.b32 {%rs271, %rs279}, %r253; // begin inline asm { cvt.f32.f16 %f286, %rs271;} // end inline asm fma.rn.ftz.f32 %f611, %f606, %f286, %f600; mov.b32 {%rs272, %rs280}, %r261; // begin inline asm { cvt.f32.f16 %f287, %rs272;} // end inline asm fma.rn.ftz.f32 %f612, %f606, %f287, %f601; mov.b32 {%rs273, %rs281}, %r269; // begin inline asm { cvt.f32.f16 %f288, %rs273;} // end inline asm fma.rn.ftz.f32 %f613, %f606, %f288, %f602; mov.b32 {%rs274, %rs282}, %r277; // begin inline asm { cvt.f32.f16 %f289, %rs274;} // end inline asm fma.rn.ftz.f32 %f614, %f606, %f289, %f603; cvt.rn.f32.s16 %f615, %rs49; sub.ftz.f32 %f616, %f615, %f10; mul.ftz.f32 %f617, %f105, %f616; // begin inline asm { cvt.f32.f16 %f290, %rs275;} // end inline asm fma.rn.ftz.f32 %f618, %f617, %f290, %f607; // begin inline asm { cvt.f32.f16 %f291, %rs276;} // end inline asm fma.rn.ftz.f32 %f619, %f617, %f291, %f608; // begin inline asm { cvt.f32.f16 %f292, %rs277;} // end inline asm fma.rn.ftz.f32 %f620, %f617, %f292, %f609; // begin inline asm { cvt.f32.f16 %f293, %rs278;} // end inline asm fma.rn.ftz.f32 %f621, %f617, %f293, %f610; // begin inline asm { cvt.f32.f16 %f294, %rs279;} // end inline asm fma.rn.ftz.f32 %f622, %f617, %f294, %f611; // begin inline asm { cvt.f32.f16 %f295, %rs280;} // end inline asm fma.rn.ftz.f32 %f623, %f617, %f295, %f612; // begin inline asm { cvt.f32.f16 %f296, %rs281;} // end inline asm fma.rn.ftz.f32 %f624, %f617, %f296, %f613; // begin inline asm { cvt.f32.f16 %f297, %rs282;} // end inline asm fma.rn.ftz.f32 %f625, %f617, %f297, %f614; ld.global.v4.u32 {%r282, %r283, %r284, %r285}, [%rd7+48]; ld.global.v4.u32 {%r290, %r291, %r292, %r293}, [%rd9+32]; ld.global.v4.u32 {%r298, %r299, %r300, %r301}, [%rd8+48]; ld.global.v4.u32 {%r306, %r307, %r308, %r309}, [%rd10+32]; ld.global.v4.u32 {%r314, %r315, %r316, %r317}, [%rd11+32]; ld.global.v4.u32 {%r322, %r323, %r324, %r325}, [%rd12+32]; ld.global.v4.u32 {%r330, %r331, %r332, %r333}, [%rd13+32]; ld.global.v4.u32 {%r338, %r339, %r340, %r341}, [%rd14+32]; cvt.rn.f32.s16 %f626, %rs51; sub.ftz.f32 %f627, %f626, %f10; mul.ftz.f32 %f628, %f105, %f627; mov.b32 {%rs283, %rs291}, %r282; // begin inline asm { cvt.f32.f16 %f298, %rs283;} // end inline asm fma.rn.ftz.f32 %f629, %f628, %f298, %f618; mov.b32 {%rs284, %rs292}, %r290; // begin inline asm { cvt.f32.f16 %f299, %rs284;} // end inline asm fma.rn.ftz.f32 %f630, %f628, %f299, %f619; mov.b32 {%rs285, %rs293}, %r298; // begin inline asm { cvt.f32.f16 %f300, %rs285;} // end inline asm fma.rn.ftz.f32 %f631, %f628, %f300, %f620; mov.b32 {%rs286, %rs294}, %r306; // begin inline asm { cvt.f32.f16 %f301, %rs286;} // end inline asm fma.rn.ftz.f32 %f632, %f628, %f301, %f621; mov.b32 {%rs287, %rs295}, %r314; // begin inline asm { cvt.f32.f16 %f302, %rs287;} // end inline asm fma.rn.ftz.f32 %f633, %f628, %f302, %f622; mov.b32 {%rs288, %rs296}, %r322; // begin inline asm { cvt.f32.f16 %f303, %rs288;} // end inline asm fma.rn.ftz.f32 %f634, %f628, %f303, %f623; mov.b32 {%rs289, %rs297}, %r330; // begin inline asm { cvt.f32.f16 %f304, %rs289;} // end inline asm fma.rn.ftz.f32 %f635, %f628, %f304, %f624; mov.b32 {%rs290, %rs298}, %r338; // begin inline asm { cvt.f32.f16 %f305, %rs290;} // end inline asm fma.rn.ftz.f32 %f636, %f628, %f305, %f625; cvt.rn.f32.s16 %f637, %rs53; sub.ftz.f32 %f638, %f637, %f10; mul.ftz.f32 %f639, %f105, %f638; // begin inline asm { cvt.f32.f16 %f306, %rs291;} // end inline asm fma.rn.ftz.f32 %f640, %f639, %f306, %f629; // begin inline asm { cvt.f32.f16 %f307, %rs292;} // end inline asm fma.rn.ftz.f32 %f641, %f639, %f307, %f630; // begin inline asm { cvt.f32.f16 %f308, %rs293;} // end inline asm fma.rn.ftz.f32 %f642, %f639, %f308, %f631; // begin inline asm { cvt.f32.f16 %f309, %rs294;} // end inline asm fma.rn.ftz.f32 %f643, %f639, %f309, %f632; // begin inline asm { cvt.f32.f16 %f310, %rs295;} // end inline asm fma.rn.ftz.f32 %f644, %f639, %f310, %f633; // begin inline asm { cvt.f32.f16 %f311, %rs296;} // end inline asm fma.rn.ftz.f32 %f645, %f639, %f311, %f634; // begin inline asm { cvt.f32.f16 %f312, %rs297;} // end inline asm fma.rn.ftz.f32 %f646, %f639, %f312, %f635; // begin inline asm { cvt.f32.f16 %f313, %rs298;} // end inline asm fma.rn.ftz.f32 %f647, %f639, %f313, %f636; cvt.rn.f32.s16 %f648, %rs55; sub.ftz.f32 %f649, %f648, %f10; mul.ftz.f32 %f650, %f105, %f649; mov.b32 {%rs299, %rs307}, %r283; // begin inline asm { cvt.f32.f16 %f314, %rs299;} // end inline asm fma.rn.ftz.f32 %f651, %f650, %f314, %f640; mov.b32 {%rs300, %rs308}, %r291; // begin inline asm { cvt.f32.f16 %f315, %rs300;} // end inline asm fma.rn.ftz.f32 %f652, %f650, %f315, %f641; mov.b32 {%rs301, %rs309}, %r299; // begin inline asm { cvt.f32.f16 %f316, %rs301;} // end inline asm fma.rn.ftz.f32 %f653, %f650, %f316, %f642; mov.b32 {%rs302, %rs310}, %r307; // begin inline asm { cvt.f32.f16 %f317, %rs302;} // end inline asm fma.rn.ftz.f32 %f654, %f650, %f317, %f643; mov.b32 {%rs303, %rs311}, %r315; // begin inline asm { cvt.f32.f16 %f318, %rs303;} // end inline asm fma.rn.ftz.f32 %f655, %f650, %f318, %f644; mov.b32 {%rs304, %rs312}, %r323; // begin inline asm { cvt.f32.f16 %f319, %rs304;} // end inline asm fma.rn.ftz.f32 %f656, %f650, %f319, %f645; mov.b32 {%rs305, %rs313}, %r331; // begin inline asm { cvt.f32.f16 %f320, %rs305;} // end inline asm fma.rn.ftz.f32 %f657, %f650, %f320, %f646; mov.b32 {%rs306, %rs314}, %r339; // begin inline asm { cvt.f32.f16 %f321, %rs306;} // end inline asm fma.rn.ftz.f32 %f658, %f650, %f321, %f647; cvt.rn.f32.s16 %f659, %rs57; sub.ftz.f32 %f660, %f659, %f10; mul.ftz.f32 %f661, %f105, %f660; // begin inline asm { cvt.f32.f16 %f322, %rs307;} // end inline asm fma.rn.ftz.f32 %f662, %f661, %f322, %f651; // begin inline asm { cvt.f32.f16 %f323, %rs308;} // end inline asm fma.rn.ftz.f32 %f663, %f661, %f323, %f652; // begin inline asm { cvt.f32.f16 %f324, %rs309;} // end inline asm fma.rn.ftz.f32 %f664, %f661, %f324, %f653; // begin inline asm { cvt.f32.f16 %f325, %rs310;} // end inline asm fma.rn.ftz.f32 %f665, %f661, %f325, %f654; // begin inline asm { cvt.f32.f16 %f326, %rs311;} // end inline asm fma.rn.ftz.f32 %f666, %f661, %f326, %f655; // begin inline asm { cvt.f32.f16 %f327, %rs312;} // end inline asm fma.rn.ftz.f32 %f667, %f661, %f327, %f656; // begin inline asm { cvt.f32.f16 %f328, %rs313;} // end inline asm fma.rn.ftz.f32 %f668, %f661, %f328, %f657; // begin inline asm { cvt.f32.f16 %f329, %rs314;} // end inline asm fma.rn.ftz.f32 %f669, %f661, %f329, %f658; cvt.rn.f32.s16 %f670, %rs59; sub.ftz.f32 %f671, %f670, %f10; mul.ftz.f32 %f672, %f105, %f671; mov.b32 {%rs315, %rs323}, %r284; // begin inline asm { cvt.f32.f16 %f330, %rs315;} // end inline asm fma.rn.ftz.f32 %f673, %f672, %f330, %f662; mov.b32 {%rs316, %rs324}, %r292; // begin inline asm { cvt.f32.f16 %f331, %rs316;} // end inline asm fma.rn.ftz.f32 %f674, %f672, %f331, %f663; mov.b32 {%rs317, %rs325}, %r300; // begin inline asm { cvt.f32.f16 %f332, %rs317;} // end inline asm fma.rn.ftz.f32 %f675, %f672, %f332, %f664; mov.b32 {%rs318, %rs326}, %r308; // begin inline asm { cvt.f32.f16 %f333, %rs318;} // end inline asm fma.rn.ftz.f32 %f676, %f672, %f333, %f665; mov.b32 {%rs319, %rs327}, %r316; // begin inline asm { cvt.f32.f16 %f334, %rs319;} // end inline asm fma.rn.ftz.f32 %f677, %f672, %f334, %f666; mov.b32 {%rs320, %rs328}, %r324; // begin inline asm { cvt.f32.f16 %f335, %rs320;} // end inline asm fma.rn.ftz.f32 %f678, %f672, %f335, %f667; mov.b32 {%rs321, %rs329}, %r332; // begin inline asm { cvt.f32.f16 %f336, %rs321;} // end inline asm fma.rn.ftz.f32 %f679, %f672, %f336, %f668; mov.b32 {%rs322, %rs330}, %r340; // begin inline asm { cvt.f32.f16 %f337, %rs322;} // end inline asm fma.rn.ftz.f32 %f680, %f672, %f337, %f669; cvt.rn.f32.s16 %f681, %rs61; sub.ftz.f32 %f682, %f681, %f10; mul.ftz.f32 %f683, %f105, %f682; // begin inline asm { cvt.f32.f16 %f338, %rs323;} // end inline asm fma.rn.ftz.f32 %f684, %f683, %f338, %f673; // begin inline asm { cvt.f32.f16 %f339, %rs324;} // end inline asm fma.rn.ftz.f32 %f685, %f683, %f339, %f674; // begin inline asm { cvt.f32.f16 %f340, %rs325;} // end inline asm fma.rn.ftz.f32 %f686, %f683, %f340, %f675; // begin inline asm { cvt.f32.f16 %f341, %rs326;} // end inline asm fma.rn.ftz.f32 %f687, %f683, %f341, %f676; // begin inline asm { cvt.f32.f16 %f342, %rs327;} // end inline asm fma.rn.ftz.f32 %f688, %f683, %f342, %f677; // begin inline asm { cvt.f32.f16 %f343, %rs328;} // end inline asm fma.rn.ftz.f32 %f689, %f683, %f343, %f678; // begin inline asm { cvt.f32.f16 %f344, %rs329;} // end inline asm fma.rn.ftz.f32 %f690, %f683, %f344, %f679; // begin inline asm { cvt.f32.f16 %f345, %rs330;} // end inline asm fma.rn.ftz.f32 %f691, %f683, %f345, %f680; cvt.rn.f32.s16 %f692, %rs63; sub.ftz.f32 %f693, %f692, %f10; mul.ftz.f32 %f694, %f105, %f693; mov.b32 {%rs331, %rs339}, %r285; // begin inline asm { cvt.f32.f16 %f346, %rs331;} // end inline asm fma.rn.ftz.f32 %f695, %f694, %f346, %f684; mov.b32 {%rs332, %rs340}, %r293; // begin inline asm { cvt.f32.f16 %f347, %rs332;} // end inline asm fma.rn.ftz.f32 %f696, %f694, %f347, %f685; mov.b32 {%rs333, %rs341}, %r301; // begin inline asm { cvt.f32.f16 %f348, %rs333;} // end inline asm fma.rn.ftz.f32 %f697, %f694, %f348, %f686; mov.b32 {%rs334, %rs342}, %r309; // begin inline asm { cvt.f32.f16 %f349, %rs334;} // end inline asm fma.rn.ftz.f32 %f698, %f694, %f349, %f687; mov.b32 {%rs335, %rs343}, %r317; // begin inline asm { cvt.f32.f16 %f350, %rs335;} // end inline asm fma.rn.ftz.f32 %f699, %f694, %f350, %f688; mov.b32 {%rs336, %rs344}, %r325; // begin inline asm { cvt.f32.f16 %f351, %rs336;} // end inline asm fma.rn.ftz.f32 %f700, %f694, %f351, %f689; mov.b32 {%rs337, %rs345}, %r333; // begin inline asm { cvt.f32.f16 %f352, %rs337;} // end inline asm fma.rn.ftz.f32 %f701, %f694, %f352, %f690; mov.b32 {%rs338, %rs346}, %r341; // begin inline asm { cvt.f32.f16 %f353, %rs338;} // end inline asm fma.rn.ftz.f32 %f702, %f694, %f353, %f691; cvt.rn.f32.s16 %f703, %rs64; sub.ftz.f32 %f704, %f703, %f10; mul.ftz.f32 %f705, %f105, %f704; // begin inline asm { cvt.f32.f16 %f354, %rs339;} // end inline asm fma.rn.ftz.f32 %f1497, %f705, %f354, %f695; // begin inline asm { cvt.f32.f16 %f355, %rs340;} // end inline asm fma.rn.ftz.f32 %f1496, %f705, %f355, %f696; // begin inline asm { cvt.f32.f16 %f356, %rs341;} // end inline asm fma.rn.ftz.f32 %f1495, %f705, %f356, %f697; // begin inline asm { cvt.f32.f16 %f357, %rs342;} // end inline asm fma.rn.ftz.f32 %f1494, %f705, %f357, %f698; // begin inline asm { cvt.f32.f16 %f358, %rs343;} // end inline asm fma.rn.ftz.f32 %f1493, %f705, %f358, %f699; // begin inline asm { cvt.f32.f16 %f359, %rs344;} // end inline asm fma.rn.ftz.f32 %f1492, %f705, %f359, %f700; // begin inline asm { cvt.f32.f16 %f360, %rs345;} // end inline asm fma.rn.ftz.f32 %f1491, %f705, %f360, %f701; // begin inline asm { cvt.f32.f16 %f361, %rs346;} // end inline asm fma.rn.ftz.f32 %f1490, %f705, %f361, %f702; bra.uni $L__BB0_8; $L__BB0_7: shl.b16 %rs603, %rs5, 4; cvt.s16.s8 %rs604, %rs603; shr.s16 %rs605, %rs604, 7; and.b16 %rs606, %rs605, -16; or.b16 %rs607, %rs606, %rs6; cvt.rn.f32.s16 %f962, %rs607; sub.ftz.f32 %f963, %f962, %f10; mul.ftz.f32 %f964, %f105, %f963; mov.b32 {%rs347, %rs355}, %r85; // begin inline asm { cvt.f32.f16 %f706, %rs347;} // end inline asm fma.rn.ftz.f32 %f965, %f964, %f706, %f1497; mov.b32 {%rs348, %rs356}, %r89; // begin inline asm { cvt.f32.f16 %f707, %rs348;} // end inline asm fma.rn.ftz.f32 %f966, %f964, %f707, %f1496; mov.b32 {%rs349, %rs357}, %r95; // begin inline asm { cvt.f32.f16 %f708, %rs349;} // end inline asm fma.rn.ftz.f32 %f967, %f964, %f708, %f1495; mov.b32 {%rs350, %rs358}, %r99; // begin inline asm { cvt.f32.f16 %f709, %rs350;} // end inline asm fma.rn.ftz.f32 %f968, %f964, %f709, %f1494; mov.b32 {%rs351, %rs359}, %r103; // begin inline asm { cvt.f32.f16 %f710, %rs351;} // end inline asm fma.rn.ftz.f32 %f969, %f964, %f710, %f1493; mov.b32 {%rs352, %rs360}, %r107; // begin inline asm { cvt.f32.f16 %f711, %rs352;} // end inline asm fma.rn.ftz.f32 %f970, %f964, %f711, %f1492; mov.b32 {%rs353, %rs361}, %r111; // begin inline asm { cvt.f32.f16 %f712, %rs353;} // end inline asm fma.rn.ftz.f32 %f971, %f964, %f712, %f1491; mov.b32 {%rs354, %rs362}, %r115; // begin inline asm { cvt.f32.f16 %f713, %rs354;} // end inline asm fma.rn.ftz.f32 %f972, %f964, %f713, %f1490; shl.b16 %rs608, %rs7, 4; cvt.s16.s8 %rs609, %rs608; shr.s16 %rs610, %rs609, 7; and.b16 %rs611, %rs610, -16; or.b16 %rs612, %rs611, %rs8; cvt.rn.f32.s16 %f973, %rs612; sub.ftz.f32 %f974, %f973, %f10; mul.ftz.f32 %f975, %f105, %f974; // begin inline asm { cvt.f32.f16 %f714, %rs355;} // end inline asm fma.rn.ftz.f32 %f976, %f975, %f714, %f965; // begin inline asm { cvt.f32.f16 %f715, %rs356;} // end inline asm fma.rn.ftz.f32 %f977, %f975, %f715, %f966; // begin inline asm { cvt.f32.f16 %f716, %rs357;} // end inline asm fma.rn.ftz.f32 %f978, %f975, %f716, %f967; // begin inline asm { cvt.f32.f16 %f717, %rs358;} // end inline asm fma.rn.ftz.f32 %f979, %f975, %f717, %f968; // begin inline asm { cvt.f32.f16 %f718, %rs359;} // end inline asm fma.rn.ftz.f32 %f980, %f975, %f718, %f969; // begin inline asm { cvt.f32.f16 %f719, %rs360;} // end inline asm fma.rn.ftz.f32 %f981, %f975, %f719, %f970; // begin inline asm { cvt.f32.f16 %f720, %rs361;} // end inline asm fma.rn.ftz.f32 %f982, %f975, %f720, %f971; // begin inline asm { cvt.f32.f16 %f721, %rs362;} // end inline asm fma.rn.ftz.f32 %f983, %f975, %f721, %f972; shl.b16 %rs613, %rs9, 4; cvt.s16.s8 %rs614, %rs613; shr.s16 %rs615, %rs614, 7; and.b16 %rs616, %rs615, -16; or.b16 %rs617, %rs616, %rs10; cvt.rn.f32.s16 %f984, %rs617; sub.ftz.f32 %f985, %f984, %f10; mul.ftz.f32 %f986, %f105, %f985; mov.b32 {%rs363, %rs371}, %r86; // begin inline asm { cvt.f32.f16 %f722, %rs363;} // end inline asm fma.rn.ftz.f32 %f987, %f986, %f722, %f976; mov.b32 {%rs364, %rs372}, %r90; // begin inline asm { cvt.f32.f16 %f723, %rs364;} // end inline asm fma.rn.ftz.f32 %f988, %f986, %f723, %f977; mov.b32 {%rs365, %rs373}, %r96; // begin inline asm { cvt.f32.f16 %f724, %rs365;} // end inline asm fma.rn.ftz.f32 %f989, %f986, %f724, %f978; mov.b32 {%rs366, %rs374}, %r100; // begin inline asm { cvt.f32.f16 %f725, %rs366;} // end inline asm fma.rn.ftz.f32 %f990, %f986, %f725, %f979; mov.b32 {%rs367, %rs375}, %r104; // begin inline asm { cvt.f32.f16 %f726, %rs367;} // end inline asm fma.rn.ftz.f32 %f991, %f986, %f726, %f980; mov.b32 {%rs368, %rs376}, %r108; // begin inline asm { cvt.f32.f16 %f727, %rs368;} // end inline asm fma.rn.ftz.f32 %f992, %f986, %f727, %f981; mov.b32 {%rs369, %rs377}, %r112; // begin inline asm { cvt.f32.f16 %f728, %rs369;} // end inline asm fma.rn.ftz.f32 %f993, %f986, %f728, %f982; mov.b32 {%rs370, %rs378}, %r116; // begin inline asm { cvt.f32.f16 %f729, %rs370;} // end inline asm fma.rn.ftz.f32 %f994, %f986, %f729, %f983; shl.b16 %rs618, %rs11, 4; cvt.s16.s8 %rs619, %rs618; shr.s16 %rs620, %rs619, 7; and.b16 %rs621, %rs620, -16; or.b16 %rs622, %rs621, %rs12; cvt.rn.f32.s16 %f995, %rs622; sub.ftz.f32 %f996, %f995, %f10; mul.ftz.f32 %f997, %f105, %f996; // begin inline asm { cvt.f32.f16 %f730, %rs371;} // end inline asm fma.rn.ftz.f32 %f998, %f997, %f730, %f987; // begin inline asm { cvt.f32.f16 %f731, %rs372;} // end inline asm fma.rn.ftz.f32 %f999, %f997, %f731, %f988; // begin inline asm { cvt.f32.f16 %f732, %rs373;} // end inline asm fma.rn.ftz.f32 %f1000, %f997, %f732, %f989; // begin inline asm { cvt.f32.f16 %f733, %rs374;} // end inline asm fma.rn.ftz.f32 %f1001, %f997, %f733, %f990; // begin inline asm { cvt.f32.f16 %f734, %rs375;} // end inline asm fma.rn.ftz.f32 %f1002, %f997, %f734, %f991; // begin inline asm { cvt.f32.f16 %f735, %rs376;} // end inline asm fma.rn.ftz.f32 %f1003, %f997, %f735, %f992; // begin inline asm { cvt.f32.f16 %f736, %rs377;} // end inline asm fma.rn.ftz.f32 %f1004, %f997, %f736, %f993; // begin inline asm { cvt.f32.f16 %f737, %rs378;} // end inline asm fma.rn.ftz.f32 %f1005, %f997, %f737, %f994; shl.b16 %rs623, %rs13, 4; cvt.s16.s8 %rs624, %rs623; shr.s16 %rs625, %rs624, 7; and.b16 %rs626, %rs625, -16; or.b16 %rs627, %rs626, %rs14; cvt.rn.f32.s16 %f1006, %rs627; sub.ftz.f32 %f1007, %f1006, %f10; mul.ftz.f32 %f1008, %f105, %f1007; mov.b32 {%rs379, %rs387}, %r87; // begin inline asm { cvt.f32.f16 %f738, %rs379;} // end inline asm fma.rn.ftz.f32 %f1009, %f1008, %f738, %f998; mov.b32 {%rs380, %rs388}, %r91; // begin inline asm { cvt.f32.f16 %f739, %rs380;} // end inline asm fma.rn.ftz.f32 %f1010, %f1008, %f739, %f999; mov.b32 {%rs381, %rs389}, %r97; // begin inline asm { cvt.f32.f16 %f740, %rs381;} // end inline asm fma.rn.ftz.f32 %f1011, %f1008, %f740, %f1000; mov.b32 {%rs382, %rs390}, %r101; // begin inline asm { cvt.f32.f16 %f741, %rs382;} // end inline asm fma.rn.ftz.f32 %f1012, %f1008, %f741, %f1001; mov.b32 {%rs383, %rs391}, %r105; // begin inline asm { cvt.f32.f16 %f742, %rs383;} // end inline asm fma.rn.ftz.f32 %f1013, %f1008, %f742, %f1002; mov.b32 {%rs384, %rs392}, %r109; // begin inline asm { cvt.f32.f16 %f743, %rs384;} // end inline asm fma.rn.ftz.f32 %f1014, %f1008, %f743, %f1003; mov.b32 {%rs385, %rs393}, %r113; // begin inline asm { cvt.f32.f16 %f744, %rs385;} // end inline asm fma.rn.ftz.f32 %f1015, %f1008, %f744, %f1004; mov.b32 {%rs386, %rs394}, %r117; // begin inline asm { cvt.f32.f16 %f745, %rs386;} // end inline asm fma.rn.ftz.f32 %f1016, %f1008, %f745, %f1005; shl.b16 %rs628, %rs15, 4; cvt.s16.s8 %rs629, %rs628; shr.s16 %rs630, %rs629, 7; and.b16 %rs631, %rs630, -16; or.b16 %rs632, %rs631, %rs16; cvt.rn.f32.s16 %f1017, %rs632; sub.ftz.f32 %f1018, %f1017, %f10; mul.ftz.f32 %f1019, %f105, %f1018; // begin inline asm { cvt.f32.f16 %f746, %rs387;} // end inline asm fma.rn.ftz.f32 %f1020, %f1019, %f746, %f1009; // begin inline asm { cvt.f32.f16 %f747, %rs388;} // end inline asm fma.rn.ftz.f32 %f1021, %f1019, %f747, %f1010; // begin inline asm { cvt.f32.f16 %f748, %rs389;} // end inline asm fma.rn.ftz.f32 %f1022, %f1019, %f748, %f1011; // begin inline asm { cvt.f32.f16 %f749, %rs390;} // end inline asm fma.rn.ftz.f32 %f1023, %f1019, %f749, %f1012; // begin inline asm { cvt.f32.f16 %f750, %rs391;} // end inline asm fma.rn.ftz.f32 %f1024, %f1019, %f750, %f1013; // begin inline asm { cvt.f32.f16 %f751, %rs392;} // end inline asm fma.rn.ftz.f32 %f1025, %f1019, %f751, %f1014; // begin inline asm { cvt.f32.f16 %f752, %rs393;} // end inline asm fma.rn.ftz.f32 %f1026, %f1019, %f752, %f1015; // begin inline asm { cvt.f32.f16 %f753, %rs394;} // end inline asm fma.rn.ftz.f32 %f1027, %f1019, %f753, %f1016; shl.b16 %rs633, %rs17, 4; cvt.s16.s8 %rs634, %rs633; shr.s16 %rs635, %rs634, 7; and.b16 %rs636, %rs635, -16; or.b16 %rs637, %rs636, %rs18; cvt.rn.f32.s16 %f1028, %rs637; sub.ftz.f32 %f1029, %f1028, %f10; mul.ftz.f32 %f1030, %f105, %f1029; mov.b32 {%rs395, %rs403}, %r88; // begin inline asm { cvt.f32.f16 %f754, %rs395;} // end inline asm fma.rn.ftz.f32 %f1031, %f1030, %f754, %f1020; mov.b32 {%rs396, %rs404}, %r92; // begin inline asm { cvt.f32.f16 %f755, %rs396;} // end inline asm fma.rn.ftz.f32 %f1032, %f1030, %f755, %f1021; mov.b32 {%rs397, %rs405}, %r98; // begin inline asm { cvt.f32.f16 %f756, %rs397;} // end inline asm fma.rn.ftz.f32 %f1033, %f1030, %f756, %f1022; mov.b32 {%rs398, %rs406}, %r102; // begin inline asm { cvt.f32.f16 %f757, %rs398;} // end inline asm fma.rn.ftz.f32 %f1034, %f1030, %f757, %f1023; mov.b32 {%rs399, %rs407}, %r106; // begin inline asm { cvt.f32.f16 %f758, %rs399;} // end inline asm fma.rn.ftz.f32 %f1035, %f1030, %f758, %f1024; mov.b32 {%rs400, %rs408}, %r110; // begin inline asm { cvt.f32.f16 %f759, %rs400;} // end inline asm fma.rn.ftz.f32 %f1036, %f1030, %f759, %f1025; mov.b32 {%rs401, %rs409}, %r114; // begin inline asm { cvt.f32.f16 %f760, %rs401;} // end inline asm fma.rn.ftz.f32 %f1037, %f1030, %f760, %f1026; mov.b32 {%rs402, %rs410}, %r118; // begin inline asm { cvt.f32.f16 %f761, %rs402;} // end inline asm fma.rn.ftz.f32 %f1038, %f1030, %f761, %f1027; shl.b16 %rs638, %rs19, 4; cvt.s16.s8 %rs639, %rs638; shr.s16 %rs640, %rs639, 7; and.b16 %rs641, %rs640, -16; or.b16 %rs642, %rs641, %rs19; cvt.rn.f32.s16 %f1039, %rs642; sub.ftz.f32 %f1040, %f1039, %f10; mul.ftz.f32 %f1041, %f105, %f1040; // begin inline asm { cvt.f32.f16 %f762, %rs403;} // end inline asm fma.rn.ftz.f32 %f1042, %f1041, %f762, %f1031; // begin inline asm { cvt.f32.f16 %f763, %rs404;} // end inline asm fma.rn.ftz.f32 %f1043, %f1041, %f763, %f1032; // begin inline asm { cvt.f32.f16 %f764, %rs405;} // end inline asm fma.rn.ftz.f32 %f1044, %f1041, %f764, %f1033; // begin inline asm { cvt.f32.f16 %f765, %rs406;} // end inline asm fma.rn.ftz.f32 %f1045, %f1041, %f765, %f1034; // begin inline asm { cvt.f32.f16 %f766, %rs407;} // end inline asm fma.rn.ftz.f32 %f1046, %f1041, %f766, %f1035; // begin inline asm { cvt.f32.f16 %f767, %rs408;} // end inline asm fma.rn.ftz.f32 %f1047, %f1041, %f767, %f1036; // begin inline asm { cvt.f32.f16 %f768, %rs409;} // end inline asm fma.rn.ftz.f32 %f1048, %f1041, %f768, %f1037; // begin inline asm { cvt.f32.f16 %f769, %rs410;} // end inline asm fma.rn.ftz.f32 %f1049, %f1041, %f769, %f1038; ld.global.v4.u32 {%r346, %r347, %r348, %r349}, [%rd7+16]; ld.global.v4.u32 {%r354, %r355, %r356, %r357}, [%rd9]; ld.global.v4.u32 {%r362, %r363, %r364, %r365}, [%rd8+16]; ld.global.v4.u32 {%r370, %r371, %r372, %r373}, [%rd10]; ld.global.v4.u32 {%r378, %r379, %r380, %r381}, [%rd11]; ld.global.v4.u32 {%r386, %r387, %r388, %r389}, [%rd12]; ld.global.v4.u32 {%r394, %r395, %r396, %r397}, [%rd13]; ld.global.v4.u32 {%r402, %r403, %r404, %r405}, [%rd14]; shl.b16 %rs643, %rs20, 4; cvt.s16.s8 %rs644, %rs643; shr.s16 %rs645, %rs644, 7; and.b16 %rs646, %rs645, -16; or.b16 %rs647, %rs646, %rs21; cvt.rn.f32.s16 %f1050, %rs647; sub.ftz.f32 %f1051, %f1050, %f10; mul.ftz.f32 %f1052, %f105, %f1051; mov.b32 {%rs411, %rs419}, %r346; // begin inline asm { cvt.f32.f16 %f770, %rs411;} // end inline asm fma.rn.ftz.f32 %f1053, %f1052, %f770, %f1042; mov.b32 {%rs412, %rs420}, %r354; // begin inline asm { cvt.f32.f16 %f771, %rs412;} // end inline asm fma.rn.ftz.f32 %f1054, %f1052, %f771, %f1043; mov.b32 {%rs413, %rs421}, %r362; // begin inline asm { cvt.f32.f16 %f772, %rs413;} // end inline asm fma.rn.ftz.f32 %f1055, %f1052, %f772, %f1044; mov.b32 {%rs414, %rs422}, %r370; // begin inline asm { cvt.f32.f16 %f773, %rs414;} // end inline asm fma.rn.ftz.f32 %f1056, %f1052, %f773, %f1045; mov.b32 {%rs415, %rs423}, %r378; // begin inline asm { cvt.f32.f16 %f774, %rs415;} // end inline asm fma.rn.ftz.f32 %f1057, %f1052, %f774, %f1046; mov.b32 {%rs416, %rs424}, %r386; // begin inline asm { cvt.f32.f16 %f775, %rs416;} // end inline asm fma.rn.ftz.f32 %f1058, %f1052, %f775, %f1047; mov.b32 {%rs417, %rs425}, %r394; // begin inline asm { cvt.f32.f16 %f776, %rs417;} // end inline asm fma.rn.ftz.f32 %f1059, %f1052, %f776, %f1048; mov.b32 {%rs418, %rs426}, %r402; // begin inline asm { cvt.f32.f16 %f777, %rs418;} // end inline asm fma.rn.ftz.f32 %f1060, %f1052, %f777, %f1049; shl.b16 %rs648, %rs22, 4; cvt.s16.s8 %rs649, %rs648; shr.s16 %rs650, %rs649, 7; and.b16 %rs651, %rs650, -16; or.b16 %rs652, %rs651, %rs23; cvt.rn.f32.s16 %f1061, %rs652; sub.ftz.f32 %f1062, %f1061, %f10; mul.ftz.f32 %f1063, %f105, %f1062; // begin inline asm { cvt.f32.f16 %f778, %rs419;} // end inline asm fma.rn.ftz.f32 %f1064, %f1063, %f778, %f1053; // begin inline asm { cvt.f32.f16 %f779, %rs420;} // end inline asm fma.rn.ftz.f32 %f1065, %f1063, %f779, %f1054; // begin inline asm { cvt.f32.f16 %f780, %rs421;} // end inline asm fma.rn.ftz.f32 %f1066, %f1063, %f780, %f1055; // begin inline asm { cvt.f32.f16 %f781, %rs422;} // end inline asm fma.rn.ftz.f32 %f1067, %f1063, %f781, %f1056; // begin inline asm { cvt.f32.f16 %f782, %rs423;} // end inline asm fma.rn.ftz.f32 %f1068, %f1063, %f782, %f1057; // begin inline asm { cvt.f32.f16 %f783, %rs424;} // end inline asm fma.rn.ftz.f32 %f1069, %f1063, %f783, %f1058; // begin inline asm { cvt.f32.f16 %f784, %rs425;} // end inline asm fma.rn.ftz.f32 %f1070, %f1063, %f784, %f1059; // begin inline asm { cvt.f32.f16 %f785, %rs426;} // end inline asm fma.rn.ftz.f32 %f1071, %f1063, %f785, %f1060; shl.b16 %rs653, %rs24, 4; cvt.s16.s8 %rs654, %rs653; shr.s16 %rs655, %rs654, 7; and.b16 %rs656, %rs655, -16; or.b16 %rs657, %rs656, %rs25; cvt.rn.f32.s16 %f1072, %rs657; sub.ftz.f32 %f1073, %f1072, %f10; mul.ftz.f32 %f1074, %f105, %f1073; mov.b32 {%rs427, %rs435}, %r347; // begin inline asm { cvt.f32.f16 %f786, %rs427;} // end inline asm fma.rn.ftz.f32 %f1075, %f1074, %f786, %f1064; mov.b32 {%rs428, %rs436}, %r355; // begin inline asm { cvt.f32.f16 %f787, %rs428;} // end inline asm fma.rn.ftz.f32 %f1076, %f1074, %f787, %f1065; mov.b32 {%rs429, %rs437}, %r363; // begin inline asm { cvt.f32.f16 %f788, %rs429;} // end inline asm fma.rn.ftz.f32 %f1077, %f1074, %f788, %f1066; mov.b32 {%rs430, %rs438}, %r371; // begin inline asm { cvt.f32.f16 %f789, %rs430;} // end inline asm fma.rn.ftz.f32 %f1078, %f1074, %f789, %f1067; mov.b32 {%rs431, %rs439}, %r379; // begin inline asm { cvt.f32.f16 %f790, %rs431;} // end inline asm fma.rn.ftz.f32 %f1079, %f1074, %f790, %f1068; mov.b32 {%rs432, %rs440}, %r387; // begin inline asm { cvt.f32.f16 %f791, %rs432;} // end inline asm fma.rn.ftz.f32 %f1080, %f1074, %f791, %f1069; mov.b32 {%rs433, %rs441}, %r395; // begin inline asm { cvt.f32.f16 %f792, %rs433;} // end inline asm fma.rn.ftz.f32 %f1081, %f1074, %f792, %f1070; mov.b32 {%rs434, %rs442}, %r403; // begin inline asm { cvt.f32.f16 %f793, %rs434;} // end inline asm fma.rn.ftz.f32 %f1082, %f1074, %f793, %f1071; shl.b16 %rs658, %rs26, 4; cvt.s16.s8 %rs659, %rs658; shr.s16 %rs660, %rs659, 7; and.b16 %rs661, %rs660, -16; or.b16 %rs662, %rs661, %rs27; cvt.rn.f32.s16 %f1083, %rs662; sub.ftz.f32 %f1084, %f1083, %f10; mul.ftz.f32 %f1085, %f105, %f1084; // begin inline asm { cvt.f32.f16 %f794, %rs435;} // end inline asm fma.rn.ftz.f32 %f1086, %f1085, %f794, %f1075; // begin inline asm { cvt.f32.f16 %f795, %rs436;} // end inline asm fma.rn.ftz.f32 %f1087, %f1085, %f795, %f1076; // begin inline asm { cvt.f32.f16 %f796, %rs437;} // end inline asm fma.rn.ftz.f32 %f1088, %f1085, %f796, %f1077; // begin inline asm { cvt.f32.f16 %f797, %rs438;} // end inline asm fma.rn.ftz.f32 %f1089, %f1085, %f797, %f1078; // begin inline asm { cvt.f32.f16 %f798, %rs439;} // end inline asm fma.rn.ftz.f32 %f1090, %f1085, %f798, %f1079; // begin inline asm { cvt.f32.f16 %f799, %rs440;} // end inline asm fma.rn.ftz.f32 %f1091, %f1085, %f799, %f1080; // begin inline asm { cvt.f32.f16 %f800, %rs441;} // end inline asm fma.rn.ftz.f32 %f1092, %f1085, %f800, %f1081; // begin inline asm { cvt.f32.f16 %f801, %rs442;} // end inline asm fma.rn.ftz.f32 %f1093, %f1085, %f801, %f1082; shl.b16 %rs663, %rs28, 4; cvt.s16.s8 %rs664, %rs663; shr.s16 %rs665, %rs664, 7; and.b16 %rs666, %rs665, -16; or.b16 %rs667, %rs666, %rs29; cvt.rn.f32.s16 %f1094, %rs667; sub.ftz.f32 %f1095, %f1094, %f10; mul.ftz.f32 %f1096, %f105, %f1095; mov.b32 {%rs443, %rs451}, %r348; // begin inline asm { cvt.f32.f16 %f802, %rs443;} // end inline asm fma.rn.ftz.f32 %f1097, %f1096, %f802, %f1086; mov.b32 {%rs444, %rs452}, %r356; // begin inline asm { cvt.f32.f16 %f803, %rs444;} // end inline asm fma.rn.ftz.f32 %f1098, %f1096, %f803, %f1087; mov.b32 {%rs445, %rs453}, %r364; // begin inline asm { cvt.f32.f16 %f804, %rs445;} // end inline asm fma.rn.ftz.f32 %f1099, %f1096, %f804, %f1088; mov.b32 {%rs446, %rs454}, %r372; // begin inline asm { cvt.f32.f16 %f805, %rs446;} // end inline asm fma.rn.ftz.f32 %f1100, %f1096, %f805, %f1089; mov.b32 {%rs447, %rs455}, %r380; // begin inline asm { cvt.f32.f16 %f806, %rs447;} // end inline asm fma.rn.ftz.f32 %f1101, %f1096, %f806, %f1090; mov.b32 {%rs448, %rs456}, %r388; // begin inline asm { cvt.f32.f16 %f807, %rs448;} // end inline asm fma.rn.ftz.f32 %f1102, %f1096, %f807, %f1091; mov.b32 {%rs449, %rs457}, %r396; // begin inline asm { cvt.f32.f16 %f808, %rs449;} // end inline asm fma.rn.ftz.f32 %f1103, %f1096, %f808, %f1092; mov.b32 {%rs450, %rs458}, %r404; // begin inline asm { cvt.f32.f16 %f809, %rs450;} // end inline asm fma.rn.ftz.f32 %f1104, %f1096, %f809, %f1093; shl.b16 %rs668, %rs30, 4; cvt.s16.s8 %rs669, %rs668; shr.s16 %rs670, %rs669, 7; and.b16 %rs671, %rs670, -16; or.b16 %rs672, %rs671, %rs31; cvt.rn.f32.s16 %f1105, %rs672; sub.ftz.f32 %f1106, %f1105, %f10; mul.ftz.f32 %f1107, %f105, %f1106; // begin inline asm { cvt.f32.f16 %f810, %rs451;} // end inline asm fma.rn.ftz.f32 %f1108, %f1107, %f810, %f1097; // begin inline asm { cvt.f32.f16 %f811, %rs452;} // end inline asm fma.rn.ftz.f32 %f1109, %f1107, %f811, %f1098; // begin inline asm { cvt.f32.f16 %f812, %rs453;} // end inline asm fma.rn.ftz.f32 %f1110, %f1107, %f812, %f1099; // begin inline asm { cvt.f32.f16 %f813, %rs454;} // end inline asm fma.rn.ftz.f32 %f1111, %f1107, %f813, %f1100; // begin inline asm { cvt.f32.f16 %f814, %rs455;} // end inline asm fma.rn.ftz.f32 %f1112, %f1107, %f814, %f1101; // begin inline asm { cvt.f32.f16 %f815, %rs456;} // end inline asm fma.rn.ftz.f32 %f1113, %f1107, %f815, %f1102; // begin inline asm { cvt.f32.f16 %f816, %rs457;} // end inline asm fma.rn.ftz.f32 %f1114, %f1107, %f816, %f1103; // begin inline asm { cvt.f32.f16 %f817, %rs458;} // end inline asm fma.rn.ftz.f32 %f1115, %f1107, %f817, %f1104; shl.b16 %rs673, %rs32, 4; cvt.s16.s8 %rs674, %rs673; shr.s16 %rs675, %rs674, 7; and.b16 %rs676, %rs675, -16; or.b16 %rs677, %rs676, %rs33; cvt.rn.f32.s16 %f1116, %rs677; sub.ftz.f32 %f1117, %f1116, %f10; mul.ftz.f32 %f1118, %f105, %f1117; mov.b32 {%rs459, %rs467}, %r349; // begin inline asm { cvt.f32.f16 %f818, %rs459;} // end inline asm fma.rn.ftz.f32 %f1119, %f1118, %f818, %f1108; mov.b32 {%rs460, %rs468}, %r357; // begin inline asm { cvt.f32.f16 %f819, %rs460;} // end inline asm fma.rn.ftz.f32 %f1120, %f1118, %f819, %f1109; mov.b32 {%rs461, %rs469}, %r365; // begin inline asm { cvt.f32.f16 %f820, %rs461;} // end inline asm fma.rn.ftz.f32 %f1121, %f1118, %f820, %f1110; mov.b32 {%rs462, %rs470}, %r373; // begin inline asm { cvt.f32.f16 %f821, %rs462;} // end inline asm fma.rn.ftz.f32 %f1122, %f1118, %f821, %f1111; mov.b32 {%rs463, %rs471}, %r381; // begin inline asm { cvt.f32.f16 %f822, %rs463;} // end inline asm fma.rn.ftz.f32 %f1123, %f1118, %f822, %f1112; mov.b32 {%rs464, %rs472}, %r389; // begin inline asm { cvt.f32.f16 %f823, %rs464;} // end inline asm fma.rn.ftz.f32 %f1124, %f1118, %f823, %f1113; mov.b32 {%rs465, %rs473}, %r397; // begin inline asm { cvt.f32.f16 %f824, %rs465;} // end inline asm fma.rn.ftz.f32 %f1125, %f1118, %f824, %f1114; mov.b32 {%rs466, %rs474}, %r405; // begin inline asm { cvt.f32.f16 %f825, %rs466;} // end inline asm fma.rn.ftz.f32 %f1126, %f1118, %f825, %f1115; shl.b16 %rs678, %rs34, 4; cvt.s16.s8 %rs679, %rs678; shr.s16 %rs680, %rs679, 7; and.b16 %rs681, %rs680, -16; or.b16 %rs682, %rs681, %rs34; cvt.rn.f32.s16 %f1127, %rs682; sub.ftz.f32 %f1128, %f1127, %f10; mul.ftz.f32 %f1129, %f105, %f1128; // begin inline asm { cvt.f32.f16 %f826, %rs467;} // end inline asm fma.rn.ftz.f32 %f1130, %f1129, %f826, %f1119; // begin inline asm { cvt.f32.f16 %f827, %rs468;} // end inline asm fma.rn.ftz.f32 %f1131, %f1129, %f827, %f1120; // begin inline asm { cvt.f32.f16 %f828, %rs469;} // end inline asm fma.rn.ftz.f32 %f1132, %f1129, %f828, %f1121; // begin inline asm { cvt.f32.f16 %f829, %rs470;} // end inline asm fma.rn.ftz.f32 %f1133, %f1129, %f829, %f1122; // begin inline asm { cvt.f32.f16 %f830, %rs471;} // end inline asm fma.rn.ftz.f32 %f1134, %f1129, %f830, %f1123; // begin inline asm { cvt.f32.f16 %f831, %rs472;} // end inline asm fma.rn.ftz.f32 %f1135, %f1129, %f831, %f1124; // begin inline asm { cvt.f32.f16 %f832, %rs473;} // end inline asm fma.rn.ftz.f32 %f1136, %f1129, %f832, %f1125; // begin inline asm { cvt.f32.f16 %f833, %rs474;} // end inline asm fma.rn.ftz.f32 %f1137, %f1129, %f833, %f1126; ld.global.v4.u32 {%r410, %r411, %r412, %r413}, [%rd7+32]; ld.global.v4.u32 {%r418, %r419, %r420, %r421}, [%rd9+16]; ld.global.v4.u32 {%r426, %r427, %r428, %r429}, [%rd8+32]; ld.global.v4.u32 {%r434, %r435, %r436, %r437}, [%rd10+16]; ld.global.v4.u32 {%r442, %r443, %r444, %r445}, [%rd11+16]; ld.global.v4.u32 {%r450, %r451, %r452, %r453}, [%rd12+16]; ld.global.v4.u32 {%r458, %r459, %r460, %r461}, [%rd13+16]; ld.global.v4.u32 {%r466, %r467, %r468, %r469}, [%rd14+16]; shl.b16 %rs683, %rs35, 4; cvt.s16.s8 %rs684, %rs683; shr.s16 %rs685, %rs684, 7; and.b16 %rs686, %rs685, -16; or.b16 %rs687, %rs686, %rs36; cvt.rn.f32.s16 %f1138, %rs687; sub.ftz.f32 %f1139, %f1138, %f10; mul.ftz.f32 %f1140, %f105, %f1139; mov.b32 {%rs475, %rs483}, %r410; // begin inline asm { cvt.f32.f16 %f834, %rs475;} // end inline asm fma.rn.ftz.f32 %f1141, %f1140, %f834, %f1130; mov.b32 {%rs476, %rs484}, %r418; // begin inline asm { cvt.f32.f16 %f835, %rs476;} // end inline asm fma.rn.ftz.f32 %f1142, %f1140, %f835, %f1131; mov.b32 {%rs477, %rs485}, %r426; // begin inline asm { cvt.f32.f16 %f836, %rs477;} // end inline asm fma.rn.ftz.f32 %f1143, %f1140, %f836, %f1132; mov.b32 {%rs478, %rs486}, %r434; // begin inline asm { cvt.f32.f16 %f837, %rs478;} // end inline asm fma.rn.ftz.f32 %f1144, %f1140, %f837, %f1133; mov.b32 {%rs479, %rs487}, %r442; // begin inline asm { cvt.f32.f16 %f838, %rs479;} // end inline asm fma.rn.ftz.f32 %f1145, %f1140, %f838, %f1134; mov.b32 {%rs480, %rs488}, %r450; // begin inline asm { cvt.f32.f16 %f839, %rs480;} // end inline asm fma.rn.ftz.f32 %f1146, %f1140, %f839, %f1135; mov.b32 {%rs481, %rs489}, %r458; // begin inline asm { cvt.f32.f16 %f840, %rs481;} // end inline asm fma.rn.ftz.f32 %f1147, %f1140, %f840, %f1136; mov.b32 {%rs482, %rs490}, %r466; // begin inline asm { cvt.f32.f16 %f841, %rs482;} // end inline asm fma.rn.ftz.f32 %f1148, %f1140, %f841, %f1137; shl.b16 %rs688, %rs37, 4; cvt.s16.s8 %rs689, %rs688; shr.s16 %rs690, %rs689, 7; and.b16 %rs691, %rs690, -16; or.b16 %rs692, %rs691, %rs38; cvt.rn.f32.s16 %f1149, %rs692; sub.ftz.f32 %f1150, %f1149, %f10; mul.ftz.f32 %f1151, %f105, %f1150; // begin inline asm { cvt.f32.f16 %f842, %rs483;} // end inline asm fma.rn.ftz.f32 %f1152, %f1151, %f842, %f1141; // begin inline asm { cvt.f32.f16 %f843, %rs484;} // end inline asm fma.rn.ftz.f32 %f1153, %f1151, %f843, %f1142; // begin inline asm { cvt.f32.f16 %f844, %rs485;} // end inline asm fma.rn.ftz.f32 %f1154, %f1151, %f844, %f1143; // begin inline asm { cvt.f32.f16 %f845, %rs486;} // end inline asm fma.rn.ftz.f32 %f1155, %f1151, %f845, %f1144; // begin inline asm { cvt.f32.f16 %f846, %rs487;} // end inline asm fma.rn.ftz.f32 %f1156, %f1151, %f846, %f1145; // begin inline asm { cvt.f32.f16 %f847, %rs488;} // end inline asm fma.rn.ftz.f32 %f1157, %f1151, %f847, %f1146; // begin inline asm { cvt.f32.f16 %f848, %rs489;} // end inline asm fma.rn.ftz.f32 %f1158, %f1151, %f848, %f1147; // begin inline asm { cvt.f32.f16 %f849, %rs490;} // end inline asm fma.rn.ftz.f32 %f1159, %f1151, %f849, %f1148; shl.b16 %rs693, %rs39, 4; cvt.s16.s8 %rs694, %rs693; shr.s16 %rs695, %rs694, 7; and.b16 %rs696, %rs695, -16; or.b16 %rs697, %rs696, %rs40; cvt.rn.f32.s16 %f1160, %rs697; sub.ftz.f32 %f1161, %f1160, %f10; mul.ftz.f32 %f1162, %f105, %f1161; mov.b32 {%rs491, %rs499}, %r411; // begin inline asm { cvt.f32.f16 %f850, %rs491;} // end inline asm fma.rn.ftz.f32 %f1163, %f1162, %f850, %f1152; mov.b32 {%rs492, %rs500}, %r419; // begin inline asm { cvt.f32.f16 %f851, %rs492;} // end inline asm fma.rn.ftz.f32 %f1164, %f1162, %f851, %f1153; mov.b32 {%rs493, %rs501}, %r427; // begin inline asm { cvt.f32.f16 %f852, %rs493;} // end inline asm fma.rn.ftz.f32 %f1165, %f1162, %f852, %f1154; mov.b32 {%rs494, %rs502}, %r435; // begin inline asm { cvt.f32.f16 %f853, %rs494;} // end inline asm fma.rn.ftz.f32 %f1166, %f1162, %f853, %f1155; mov.b32 {%rs495, %rs503}, %r443; // begin inline asm { cvt.f32.f16 %f854, %rs495;} // end inline asm fma.rn.ftz.f32 %f1167, %f1162, %f854, %f1156; mov.b32 {%rs496, %rs504}, %r451; // begin inline asm { cvt.f32.f16 %f855, %rs496;} // end inline asm fma.rn.ftz.f32 %f1168, %f1162, %f855, %f1157; mov.b32 {%rs497, %rs505}, %r459; // begin inline asm { cvt.f32.f16 %f856, %rs497;} // end inline asm fma.rn.ftz.f32 %f1169, %f1162, %f856, %f1158; mov.b32 {%rs498, %rs506}, %r467; // begin inline asm { cvt.f32.f16 %f857, %rs498;} // end inline asm fma.rn.ftz.f32 %f1170, %f1162, %f857, %f1159; shl.b16 %rs698, %rs41, 4; cvt.s16.s8 %rs699, %rs698; shr.s16 %rs700, %rs699, 7; and.b16 %rs701, %rs700, -16; or.b16 %rs702, %rs701, %rs42; cvt.rn.f32.s16 %f1171, %rs702; sub.ftz.f32 %f1172, %f1171, %f10; mul.ftz.f32 %f1173, %f105, %f1172; // begin inline asm { cvt.f32.f16 %f858, %rs499;} // end inline asm fma.rn.ftz.f32 %f1174, %f1173, %f858, %f1163; // begin inline asm { cvt.f32.f16 %f859, %rs500;} // end inline asm fma.rn.ftz.f32 %f1175, %f1173, %f859, %f1164; // begin inline asm { cvt.f32.f16 %f860, %rs501;} // end inline asm fma.rn.ftz.f32 %f1176, %f1173, %f860, %f1165; // begin inline asm { cvt.f32.f16 %f861, %rs502;} // end inline asm fma.rn.ftz.f32 %f1177, %f1173, %f861, %f1166; // begin inline asm { cvt.f32.f16 %f862, %rs503;} // end inline asm fma.rn.ftz.f32 %f1178, %f1173, %f862, %f1167; // begin inline asm { cvt.f32.f16 %f863, %rs504;} // end inline asm fma.rn.ftz.f32 %f1179, %f1173, %f863, %f1168; // begin inline asm { cvt.f32.f16 %f864, %rs505;} // end inline asm fma.rn.ftz.f32 %f1180, %f1173, %f864, %f1169; // begin inline asm { cvt.f32.f16 %f865, %rs506;} // end inline asm fma.rn.ftz.f32 %f1181, %f1173, %f865, %f1170; shl.b16 %rs703, %rs43, 4; cvt.s16.s8 %rs704, %rs703; shr.s16 %rs705, %rs704, 7; and.b16 %rs706, %rs705, -16; or.b16 %rs707, %rs706, %rs44; cvt.rn.f32.s16 %f1182, %rs707; sub.ftz.f32 %f1183, %f1182, %f10; mul.ftz.f32 %f1184, %f105, %f1183; mov.b32 {%rs507, %rs515}, %r412; // begin inline asm { cvt.f32.f16 %f866, %rs507;} // end inline asm fma.rn.ftz.f32 %f1185, %f1184, %f866, %f1174; mov.b32 {%rs508, %rs516}, %r420; // begin inline asm { cvt.f32.f16 %f867, %rs508;} // end inline asm fma.rn.ftz.f32 %f1186, %f1184, %f867, %f1175; mov.b32 {%rs509, %rs517}, %r428; // begin inline asm { cvt.f32.f16 %f868, %rs509;} // end inline asm fma.rn.ftz.f32 %f1187, %f1184, %f868, %f1176; mov.b32 {%rs510, %rs518}, %r436; // begin inline asm { cvt.f32.f16 %f869, %rs510;} // end inline asm fma.rn.ftz.f32 %f1188, %f1184, %f869, %f1177; mov.b32 {%rs511, %rs519}, %r444; // begin inline asm { cvt.f32.f16 %f870, %rs511;} // end inline asm fma.rn.ftz.f32 %f1189, %f1184, %f870, %f1178; mov.b32 {%rs512, %rs520}, %r452; // begin inline asm { cvt.f32.f16 %f871, %rs512;} // end inline asm fma.rn.ftz.f32 %f1190, %f1184, %f871, %f1179; mov.b32 {%rs513, %rs521}, %r460; // begin inline asm { cvt.f32.f16 %f872, %rs513;} // end inline asm fma.rn.ftz.f32 %f1191, %f1184, %f872, %f1180; mov.b32 {%rs514, %rs522}, %r468; // begin inline asm { cvt.f32.f16 %f873, %rs514;} // end inline asm fma.rn.ftz.f32 %f1192, %f1184, %f873, %f1181; shl.b16 %rs708, %rs45, 4; cvt.s16.s8 %rs709, %rs708; shr.s16 %rs710, %rs709, 7; and.b16 %rs711, %rs710, -16; or.b16 %rs712, %rs711, %rs46; cvt.rn.f32.s16 %f1193, %rs712; sub.ftz.f32 %f1194, %f1193, %f10; mul.ftz.f32 %f1195, %f105, %f1194; // begin inline asm { cvt.f32.f16 %f874, %rs515;} // end inline asm fma.rn.ftz.f32 %f1196, %f1195, %f874, %f1185; // begin inline asm { cvt.f32.f16 %f875, %rs516;} // end inline asm fma.rn.ftz.f32 %f1197, %f1195, %f875, %f1186; // begin inline asm { cvt.f32.f16 %f876, %rs517;} // end inline asm fma.rn.ftz.f32 %f1198, %f1195, %f876, %f1187; // begin inline asm { cvt.f32.f16 %f877, %rs518;} // end inline asm fma.rn.ftz.f32 %f1199, %f1195, %f877, %f1188; // begin inline asm { cvt.f32.f16 %f878, %rs519;} // end inline asm fma.rn.ftz.f32 %f1200, %f1195, %f878, %f1189; // begin inline asm { cvt.f32.f16 %f879, %rs520;} // end inline asm fma.rn.ftz.f32 %f1201, %f1195, %f879, %f1190; // begin inline asm { cvt.f32.f16 %f880, %rs521;} // end inline asm fma.rn.ftz.f32 %f1202, %f1195, %f880, %f1191; // begin inline asm { cvt.f32.f16 %f881, %rs522;} // end inline asm fma.rn.ftz.f32 %f1203, %f1195, %f881, %f1192; shl.b16 %rs713, %rs47, 4; cvt.s16.s8 %rs714, %rs713; shr.s16 %rs715, %rs714, 7; and.b16 %rs716, %rs715, -16; or.b16 %rs717, %rs716, %rs48; cvt.rn.f32.s16 %f1204, %rs717; sub.ftz.f32 %f1205, %f1204, %f10; mul.ftz.f32 %f1206, %f105, %f1205; mov.b32 {%rs523, %rs531}, %r413; // begin inline asm { cvt.f32.f16 %f882, %rs523;} // end inline asm fma.rn.ftz.f32 %f1207, %f1206, %f882, %f1196; mov.b32 {%rs524, %rs532}, %r421; // begin inline asm { cvt.f32.f16 %f883, %rs524;} // end inline asm fma.rn.ftz.f32 %f1208, %f1206, %f883, %f1197; mov.b32 {%rs525, %rs533}, %r429; // begin inline asm { cvt.f32.f16 %f884, %rs525;} // end inline asm fma.rn.ftz.f32 %f1209, %f1206, %f884, %f1198; mov.b32 {%rs526, %rs534}, %r437; // begin inline asm { cvt.f32.f16 %f885, %rs526;} // end inline asm fma.rn.ftz.f32 %f1210, %f1206, %f885, %f1199; mov.b32 {%rs527, %rs535}, %r445; // begin inline asm { cvt.f32.f16 %f886, %rs527;} // end inline asm fma.rn.ftz.f32 %f1211, %f1206, %f886, %f1200; mov.b32 {%rs528, %rs536}, %r453; // begin inline asm { cvt.f32.f16 %f887, %rs528;} // end inline asm fma.rn.ftz.f32 %f1212, %f1206, %f887, %f1201; mov.b32 {%rs529, %rs537}, %r461; // begin inline asm { cvt.f32.f16 %f888, %rs529;} // end inline asm fma.rn.ftz.f32 %f1213, %f1206, %f888, %f1202; mov.b32 {%rs530, %rs538}, %r469; // begin inline asm { cvt.f32.f16 %f889, %rs530;} // end inline asm fma.rn.ftz.f32 %f1214, %f1206, %f889, %f1203; shl.b16 %rs718, %rs49, 4; cvt.s16.s8 %rs719, %rs718; shr.s16 %rs720, %rs719, 7; and.b16 %rs721, %rs720, -16; or.b16 %rs722, %rs721, %rs49; cvt.rn.f32.s16 %f1215, %rs722; sub.ftz.f32 %f1216, %f1215, %f10; mul.ftz.f32 %f1217, %f105, %f1216; // begin inline asm { cvt.f32.f16 %f890, %rs531;} // end inline asm fma.rn.ftz.f32 %f1218, %f1217, %f890, %f1207; // begin inline asm { cvt.f32.f16 %f891, %rs532;} // end inline asm fma.rn.ftz.f32 %f1219, %f1217, %f891, %f1208; // begin inline asm { cvt.f32.f16 %f892, %rs533;} // end inline asm fma.rn.ftz.f32 %f1220, %f1217, %f892, %f1209; // begin inline asm { cvt.f32.f16 %f893, %rs534;} // end inline asm fma.rn.ftz.f32 %f1221, %f1217, %f893, %f1210; // begin inline asm { cvt.f32.f16 %f894, %rs535;} // end inline asm fma.rn.ftz.f32 %f1222, %f1217, %f894, %f1211; // begin inline asm { cvt.f32.f16 %f895, %rs536;} // end inline asm fma.rn.ftz.f32 %f1223, %f1217, %f895, %f1212; // begin inline asm { cvt.f32.f16 %f896, %rs537;} // end inline asm fma.rn.ftz.f32 %f1224, %f1217, %f896, %f1213; // begin inline asm { cvt.f32.f16 %f897, %rs538;} // end inline asm fma.rn.ftz.f32 %f1225, %f1217, %f897, %f1214; ld.global.v4.u32 {%r474, %r475, %r476, %r477}, [%rd7+48]; ld.global.v4.u32 {%r482, %r483, %r484, %r485}, [%rd9+32]; ld.global.v4.u32 {%r490, %r491, %r492, %r493}, [%rd8+48]; ld.global.v4.u32 {%r498, %r499, %r500, %r501}, [%rd10+32]; ld.global.v4.u32 {%r506, %r507, %r508, %r509}, [%rd11+32]; ld.global.v4.u32 {%r514, %r515, %r516, %r517}, [%rd12+32]; ld.global.v4.u32 {%r522, %r523, %r524, %r525}, [%rd13+32]; ld.global.v4.u32 {%r530, %r531, %r532, %r533}, [%rd14+32]; shl.b16 %rs723, %rs50, 4; cvt.s16.s8 %rs724, %rs723; shr.s16 %rs725, %rs724, 7; and.b16 %rs726, %rs725, -16; or.b16 %rs727, %rs726, %rs51; cvt.rn.f32.s16 %f1226, %rs727; sub.ftz.f32 %f1227, %f1226, %f10; mul.ftz.f32 %f1228, %f105, %f1227; mov.b32 {%rs539, %rs547}, %r474; // begin inline asm { cvt.f32.f16 %f898, %rs539;} // end inline asm fma.rn.ftz.f32 %f1229, %f1228, %f898, %f1218; mov.b32 {%rs540, %rs548}, %r482; // begin inline asm { cvt.f32.f16 %f899, %rs540;} // end inline asm fma.rn.ftz.f32 %f1230, %f1228, %f899, %f1219; mov.b32 {%rs541, %rs549}, %r490; // begin inline asm { cvt.f32.f16 %f900, %rs541;} // end inline asm fma.rn.ftz.f32 %f1231, %f1228, %f900, %f1220; mov.b32 {%rs542, %rs550}, %r498; // begin inline asm { cvt.f32.f16 %f901, %rs542;} // end inline asm fma.rn.ftz.f32 %f1232, %f1228, %f901, %f1221; mov.b32 {%rs543, %rs551}, %r506; // begin inline asm { cvt.f32.f16 %f902, %rs543;} // end inline asm fma.rn.ftz.f32 %f1233, %f1228, %f902, %f1222; mov.b32 {%rs544, %rs552}, %r514; // begin inline asm { cvt.f32.f16 %f903, %rs544;} // end inline asm fma.rn.ftz.f32 %f1234, %f1228, %f903, %f1223; mov.b32 {%rs545, %rs553}, %r522; // begin inline asm { cvt.f32.f16 %f904, %rs545;} // end inline asm fma.rn.ftz.f32 %f1235, %f1228, %f904, %f1224; mov.b32 {%rs546, %rs554}, %r530; // begin inline asm { cvt.f32.f16 %f905, %rs546;} // end inline asm fma.rn.ftz.f32 %f1236, %f1228, %f905, %f1225; shl.b16 %rs728, %rs52, 4; cvt.s16.s8 %rs729, %rs728; shr.s16 %rs730, %rs729, 7; and.b16 %rs731, %rs730, -16; or.b16 %rs732, %rs731, %rs53; cvt.rn.f32.s16 %f1237, %rs732; sub.ftz.f32 %f1238, %f1237, %f10; mul.ftz.f32 %f1239, %f105, %f1238; // begin inline asm { cvt.f32.f16 %f906, %rs547;} // end inline asm fma.rn.ftz.f32 %f1240, %f1239, %f906, %f1229; // begin inline asm { cvt.f32.f16 %f907, %rs548;} // end inline asm fma.rn.ftz.f32 %f1241, %f1239, %f907, %f1230; // begin inline asm { cvt.f32.f16 %f908, %rs549;} // end inline asm fma.rn.ftz.f32 %f1242, %f1239, %f908, %f1231; // begin inline asm { cvt.f32.f16 %f909, %rs550;} // end inline asm fma.rn.ftz.f32 %f1243, %f1239, %f909, %f1232; // begin inline asm { cvt.f32.f16 %f910, %rs551;} // end inline asm fma.rn.ftz.f32 %f1244, %f1239, %f910, %f1233; // begin inline asm { cvt.f32.f16 %f911, %rs552;} // end inline asm fma.rn.ftz.f32 %f1245, %f1239, %f911, %f1234; // begin inline asm { cvt.f32.f16 %f912, %rs553;} // end inline asm fma.rn.ftz.f32 %f1246, %f1239, %f912, %f1235; // begin inline asm { cvt.f32.f16 %f913, %rs554;} // end inline asm fma.rn.ftz.f32 %f1247, %f1239, %f913, %f1236; shl.b16 %rs733, %rs54, 4; cvt.s16.s8 %rs734, %rs733; shr.s16 %rs735, %rs734, 7; and.b16 %rs736, %rs735, -16; or.b16 %rs737, %rs736, %rs55; cvt.rn.f32.s16 %f1248, %rs737; sub.ftz.f32 %f1249, %f1248, %f10; mul.ftz.f32 %f1250, %f105, %f1249; mov.b32 {%rs555, %rs563}, %r475; // begin inline asm { cvt.f32.f16 %f914, %rs555;} // end inline asm fma.rn.ftz.f32 %f1251, %f1250, %f914, %f1240; mov.b32 {%rs556, %rs564}, %r483; // begin inline asm { cvt.f32.f16 %f915, %rs556;} // end inline asm fma.rn.ftz.f32 %f1252, %f1250, %f915, %f1241; mov.b32 {%rs557, %rs565}, %r491; // begin inline asm { cvt.f32.f16 %f916, %rs557;} // end inline asm fma.rn.ftz.f32 %f1253, %f1250, %f916, %f1242; mov.b32 {%rs558, %rs566}, %r499; // begin inline asm { cvt.f32.f16 %f917, %rs558;} // end inline asm fma.rn.ftz.f32 %f1254, %f1250, %f917, %f1243; mov.b32 {%rs559, %rs567}, %r507; // begin inline asm { cvt.f32.f16 %f918, %rs559;} // end inline asm fma.rn.ftz.f32 %f1255, %f1250, %f918, %f1244; mov.b32 {%rs560, %rs568}, %r515; // begin inline asm { cvt.f32.f16 %f919, %rs560;} // end inline asm fma.rn.ftz.f32 %f1256, %f1250, %f919, %f1245; mov.b32 {%rs561, %rs569}, %r523; // begin inline asm { cvt.f32.f16 %f920, %rs561;} // end inline asm fma.rn.ftz.f32 %f1257, %f1250, %f920, %f1246; mov.b32 {%rs562, %rs570}, %r531; // begin inline asm { cvt.f32.f16 %f921, %rs562;} // end inline asm fma.rn.ftz.f32 %f1258, %f1250, %f921, %f1247; shl.b16 %rs738, %rs56, 4; cvt.s16.s8 %rs739, %rs738; shr.s16 %rs740, %rs739, 7; and.b16 %rs741, %rs740, -16; or.b16 %rs742, %rs741, %rs57; cvt.rn.f32.s16 %f1259, %rs742; sub.ftz.f32 %f1260, %f1259, %f10; mul.ftz.f32 %f1261, %f105, %f1260; // begin inline asm { cvt.f32.f16 %f922, %rs563;} // end inline asm fma.rn.ftz.f32 %f1262, %f1261, %f922, %f1251; // begin inline asm { cvt.f32.f16 %f923, %rs564;} // end inline asm fma.rn.ftz.f32 %f1263, %f1261, %f923, %f1252; // begin inline asm { cvt.f32.f16 %f924, %rs565;} // end inline asm fma.rn.ftz.f32 %f1264, %f1261, %f924, %f1253; // begin inline asm { cvt.f32.f16 %f925, %rs566;} // end inline asm fma.rn.ftz.f32 %f1265, %f1261, %f925, %f1254; // begin inline asm { cvt.f32.f16 %f926, %rs567;} // end inline asm fma.rn.ftz.f32 %f1266, %f1261, %f926, %f1255; // begin inline asm { cvt.f32.f16 %f927, %rs568;} // end inline asm fma.rn.ftz.f32 %f1267, %f1261, %f927, %f1256; // begin inline asm { cvt.f32.f16 %f928, %rs569;} // end inline asm fma.rn.ftz.f32 %f1268, %f1261, %f928, %f1257; // begin inline asm { cvt.f32.f16 %f929, %rs570;} // end inline asm fma.rn.ftz.f32 %f1269, %f1261, %f929, %f1258; shl.b16 %rs743, %rs58, 4; cvt.s16.s8 %rs744, %rs743; shr.s16 %rs745, %rs744, 7; and.b16 %rs746, %rs745, -16; or.b16 %rs747, %rs746, %rs59; cvt.rn.f32.s16 %f1270, %rs747; sub.ftz.f32 %f1271, %f1270, %f10; mul.ftz.f32 %f1272, %f105, %f1271; mov.b32 {%rs571, %rs579}, %r476; // begin inline asm { cvt.f32.f16 %f930, %rs571;} // end inline asm fma.rn.ftz.f32 %f1273, %f1272, %f930, %f1262; mov.b32 {%rs572, %rs580}, %r484; // begin inline asm { cvt.f32.f16 %f931, %rs572;} // end inline asm fma.rn.ftz.f32 %f1274, %f1272, %f931, %f1263; mov.b32 {%rs573, %rs581}, %r492; // begin inline asm { cvt.f32.f16 %f932, %rs573;} // end inline asm fma.rn.ftz.f32 %f1275, %f1272, %f932, %f1264; mov.b32 {%rs574, %rs582}, %r500; // begin inline asm { cvt.f32.f16 %f933, %rs574;} // end inline asm fma.rn.ftz.f32 %f1276, %f1272, %f933, %f1265; mov.b32 {%rs575, %rs583}, %r508; // begin inline asm { cvt.f32.f16 %f934, %rs575;} // end inline asm fma.rn.ftz.f32 %f1277, %f1272, %f934, %f1266; mov.b32 {%rs576, %rs584}, %r516; // begin inline asm { cvt.f32.f16 %f935, %rs576;} // end inline asm fma.rn.ftz.f32 %f1278, %f1272, %f935, %f1267; mov.b32 {%rs577, %rs585}, %r524; // begin inline asm { cvt.f32.f16 %f936, %rs577;} // end inline asm fma.rn.ftz.f32 %f1279, %f1272, %f936, %f1268; mov.b32 {%rs578, %rs586}, %r532; // begin inline asm { cvt.f32.f16 %f937, %rs578;} // end inline asm fma.rn.ftz.f32 %f1280, %f1272, %f937, %f1269; shl.b16 %rs748, %rs60, 4; cvt.s16.s8 %rs749, %rs748; shr.s16 %rs750, %rs749, 7; and.b16 %rs751, %rs750, -16; or.b16 %rs752, %rs751, %rs61; cvt.rn.f32.s16 %f1281, %rs752; sub.ftz.f32 %f1282, %f1281, %f10; mul.ftz.f32 %f1283, %f105, %f1282; // begin inline asm { cvt.f32.f16 %f938, %rs579;} // end inline asm fma.rn.ftz.f32 %f1284, %f1283, %f938, %f1273; // begin inline asm { cvt.f32.f16 %f939, %rs580;} // end inline asm fma.rn.ftz.f32 %f1285, %f1283, %f939, %f1274; // begin inline asm { cvt.f32.f16 %f940, %rs581;} // end inline asm fma.rn.ftz.f32 %f1286, %f1283, %f940, %f1275; // begin inline asm { cvt.f32.f16 %f941, %rs582;} // end inline asm fma.rn.ftz.f32 %f1287, %f1283, %f941, %f1276; // begin inline asm { cvt.f32.f16 %f942, %rs583;} // end inline asm fma.rn.ftz.f32 %f1288, %f1283, %f942, %f1277; // begin inline asm { cvt.f32.f16 %f943, %rs584;} // end inline asm fma.rn.ftz.f32 %f1289, %f1283, %f943, %f1278; // begin inline asm { cvt.f32.f16 %f944, %rs585;} // end inline asm fma.rn.ftz.f32 %f1290, %f1283, %f944, %f1279; // begin inline asm { cvt.f32.f16 %f945, %rs586;} // end inline asm fma.rn.ftz.f32 %f1291, %f1283, %f945, %f1280; shl.b16 %rs753, %rs62, 4; cvt.s16.s8 %rs754, %rs753; shr.s16 %rs755, %rs754, 7; and.b16 %rs756, %rs755, -16; or.b16 %rs757, %rs756, %rs63; cvt.rn.f32.s16 %f1292, %rs757; sub.ftz.f32 %f1293, %f1292, %f10; mul.ftz.f32 %f1294, %f105, %f1293; mov.b32 {%rs587, %rs595}, %r477; // begin inline asm { cvt.f32.f16 %f946, %rs587;} // end inline asm fma.rn.ftz.f32 %f1295, %f1294, %f946, %f1284; mov.b32 {%rs588, %rs596}, %r485; // begin inline asm { cvt.f32.f16 %f947, %rs588;} // end inline asm fma.rn.ftz.f32 %f1296, %f1294, %f947, %f1285; mov.b32 {%rs589, %rs597}, %r493; // begin inline asm { cvt.f32.f16 %f948, %rs589;} // end inline asm fma.rn.ftz.f32 %f1297, %f1294, %f948, %f1286; mov.b32 {%rs590, %rs598}, %r501; // begin inline asm { cvt.f32.f16 %f949, %rs590;} // end inline asm fma.rn.ftz.f32 %f1298, %f1294, %f949, %f1287; mov.b32 {%rs591, %rs599}, %r509; // begin inline asm { cvt.f32.f16 %f950, %rs591;} // end inline asm fma.rn.ftz.f32 %f1299, %f1294, %f950, %f1288; mov.b32 {%rs592, %rs600}, %r517; // begin inline asm { cvt.f32.f16 %f951, %rs592;} // end inline asm fma.rn.ftz.f32 %f1300, %f1294, %f951, %f1289; mov.b32 {%rs593, %rs601}, %r525; // begin inline asm { cvt.f32.f16 %f952, %rs593;} // end inline asm fma.rn.ftz.f32 %f1301, %f1294, %f952, %f1290; mov.b32 {%rs594, %rs602}, %r533; // begin inline asm { cvt.f32.f16 %f953, %rs594;} // end inline asm fma.rn.ftz.f32 %f1302, %f1294, %f953, %f1291; shl.b16 %rs758, %rs64, 4; cvt.s16.s8 %rs759, %rs758; shr.s16 %rs760, %rs759, 7; and.b16 %rs761, %rs760, -16; or.b16 %rs762, %rs761, %rs64; cvt.rn.f32.s16 %f1303, %rs762; sub.ftz.f32 %f1304, %f1303, %f10; mul.ftz.f32 %f1305, %f105, %f1304; // begin inline asm { cvt.f32.f16 %f954, %rs595;} // end inline asm fma.rn.ftz.f32 %f1497, %f1305, %f954, %f1295; // begin inline asm { cvt.f32.f16 %f955, %rs596;} // end inline asm fma.rn.ftz.f32 %f1496, %f1305, %f955, %f1296; // begin inline asm { cvt.f32.f16 %f956, %rs597;} // end inline asm fma.rn.ftz.f32 %f1495, %f1305, %f956, %f1297; // begin inline asm { cvt.f32.f16 %f957, %rs598;} // end inline asm fma.rn.ftz.f32 %f1494, %f1305, %f957, %f1298; // begin inline asm { cvt.f32.f16 %f958, %rs599;} // end inline asm fma.rn.ftz.f32 %f1493, %f1305, %f958, %f1299; // begin inline asm { cvt.f32.f16 %f959, %rs600;} // end inline asm fma.rn.ftz.f32 %f1492, %f1305, %f959, %f1300; // begin inline asm { cvt.f32.f16 %f960, %rs601;} // end inline asm fma.rn.ftz.f32 %f1491, %f1305, %f960, %f1301; // begin inline asm { cvt.f32.f16 %f961, %rs602;} // end inline asm fma.rn.ftz.f32 %f1490, %f1305, %f961, %f1302; $L__BB0_8: add.s32 %r762, %r762, 4; shl.b32 %r538, %r762, 5; add.s32 %r761, %r538, %r66; shl.b32 %r760, %r761, 2; setp.lt.u32 %p7, %r760, %r63; @%p7 bra $L__BB0_2; $L__BB0_9: mov.u32 %r540, %tid.y; shl.b32 %r541, %r540, 5; add.s32 %r52, %r541, %r66; setp.lt.u32 %p8, %r52, 32; shl.b32 %r543, %r52, 2; mov.u32 %r544, _ZZ9gemv_int4ILi4ELi128ELi8EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage; add.s32 %r545, %r544, %r543; @%p8 bra $L__BB0_11; add.s32 %r752, %r545, -112; st.shared.f32 [%r752], %f1497; $L__BB0_11: setp.gt.u32 %p9, %r52, 31; bar.sync 0; mad.lo.s32 %r54, %r52, 12, %r544; @%p9 bra $L__BB0_13; mov.u32 %r564, 16; ld.shared.f32 %f1321, [%r54+16]; add.ftz.f32 %f1322, %f1497, %f1321; ld.shared.f32 %f1323, [%r54+20]; add.ftz.f32 %f1324, %f1322, %f1323; ld.shared.f32 %f1325, [%r54+24]; add.ftz.f32 %f1308, %f1324, %f1325; mov.u32 %r552, 1; mov.u32 %r565, 31; mov.u32 %r566, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1308, %r552, %r565, %r566; @p add.f32 r0, r0, %f1308; mov.f32 %f1306, r0;} // end inline asm mov.u32 %r555, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1306, %r555, %r565, %r566; @p add.f32 r0, r0, %f1306; mov.f32 %f1309, r0;} // end inline asm mov.u32 %r558, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1309, %r558, %r565, %r566; @p add.f32 r0, r0, %f1309; mov.f32 %f1312, r0;} // end inline asm mov.u32 %r561, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1312, %r561, %r565, %r566; @p add.f32 r0, r0, %f1312; mov.f32 %f1315, r0;} // end inline asm // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1315, %r564, %r565, %r566; @p add.f32 r0, r0, %f1315; mov.f32 %f1497, r0;} // end inline asm $L__BB0_13: @%p8 bra $L__BB0_15; add.s32 %r753, %r545, -112; st.shared.f32 [%r753+640], %f1496; $L__BB0_15: bar.sync 0; @%p9 bra $L__BB0_17; ld.shared.f32 %f1341, [%r54+656]; add.ftz.f32 %f1342, %f1496, %f1341; ld.shared.f32 %f1343, [%r54+660]; add.ftz.f32 %f1344, %f1342, %f1343; ld.shared.f32 %f1345, [%r54+664]; add.ftz.f32 %f1328, %f1344, %f1345; mov.u32 %r576, 1; mov.u32 %r589, 31; mov.u32 %r590, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1328, %r576, %r589, %r590; @p add.f32 r0, r0, %f1328; mov.f32 %f1326, r0;} // end inline asm mov.u32 %r579, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1326, %r579, %r589, %r590; @p add.f32 r0, r0, %f1326; mov.f32 %f1329, r0;} // end inline asm mov.u32 %r582, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1329, %r582, %r589, %r590; @p add.f32 r0, r0, %f1329; mov.f32 %f1332, r0;} // end inline asm mov.u32 %r585, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1332, %r585, %r589, %r590; @p add.f32 r0, r0, %f1332; mov.f32 %f1335, r0;} // end inline asm mov.u32 %r588, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1335, %r588, %r589, %r590; @p add.f32 r0, r0, %f1335; mov.f32 %f1496, r0;} // end inline asm $L__BB0_17: @%p8 bra $L__BB0_19; add.s32 %r754, %r545, -112; st.shared.f32 [%r754+1280], %f1495; $L__BB0_19: bar.sync 0; @%p9 bra $L__BB0_21; ld.shared.f32 %f1361, [%r54+1296]; add.ftz.f32 %f1362, %f1495, %f1361; ld.shared.f32 %f1363, [%r54+1300]; add.ftz.f32 %f1364, %f1362, %f1363; ld.shared.f32 %f1365, [%r54+1304]; add.ftz.f32 %f1348, %f1364, %f1365; mov.u32 %r600, 1; mov.u32 %r613, 31; mov.u32 %r614, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1348, %r600, %r613, %r614; @p add.f32 r0, r0, %f1348; mov.f32 %f1346, r0;} // end inline asm mov.u32 %r603, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1346, %r603, %r613, %r614; @p add.f32 r0, r0, %f1346; mov.f32 %f1349, r0;} // end inline asm mov.u32 %r606, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1349, %r606, %r613, %r614; @p add.f32 r0, r0, %f1349; mov.f32 %f1352, r0;} // end inline asm mov.u32 %r609, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1352, %r609, %r613, %r614; @p add.f32 r0, r0, %f1352; mov.f32 %f1355, r0;} // end inline asm mov.u32 %r612, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1355, %r612, %r613, %r614; @p add.f32 r0, r0, %f1355; mov.f32 %f1495, r0;} // end inline asm $L__BB0_21: @%p8 bra $L__BB0_23; add.s32 %r755, %r545, -112; st.shared.f32 [%r755+1920], %f1494; $L__BB0_23: bar.sync 0; @%p9 bra $L__BB0_25; ld.shared.f32 %f1381, [%r54+1936]; add.ftz.f32 %f1382, %f1494, %f1381; ld.shared.f32 %f1383, [%r54+1940]; add.ftz.f32 %f1384, %f1382, %f1383; ld.shared.f32 %f1385, [%r54+1944]; add.ftz.f32 %f1368, %f1384, %f1385; mov.u32 %r624, 1; mov.u32 %r637, 31; mov.u32 %r638, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1368, %r624, %r637, %r638; @p add.f32 r0, r0, %f1368; mov.f32 %f1366, r0;} // end inline asm mov.u32 %r627, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1366, %r627, %r637, %r638; @p add.f32 r0, r0, %f1366; mov.f32 %f1369, r0;} // end inline asm mov.u32 %r630, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1369, %r630, %r637, %r638; @p add.f32 r0, r0, %f1369; mov.f32 %f1372, r0;} // end inline asm mov.u32 %r633, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1372, %r633, %r637, %r638; @p add.f32 r0, r0, %f1372; mov.f32 %f1375, r0;} // end inline asm mov.u32 %r636, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1375, %r636, %r637, %r638; @p add.f32 r0, r0, %f1375; mov.f32 %f1494, r0;} // end inline asm $L__BB0_25: @%p8 bra $L__BB0_27; add.s32 %r756, %r545, -112; st.shared.f32 [%r756+2560], %f1493; $L__BB0_27: bar.sync 0; @%p9 bra $L__BB0_29; ld.shared.f32 %f1401, [%r54+2576]; add.ftz.f32 %f1402, %f1493, %f1401; ld.shared.f32 %f1403, [%r54+2580]; add.ftz.f32 %f1404, %f1402, %f1403; ld.shared.f32 %f1405, [%r54+2584]; add.ftz.f32 %f1388, %f1404, %f1405; mov.u32 %r648, 1; mov.u32 %r661, 31; mov.u32 %r662, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1388, %r648, %r661, %r662; @p add.f32 r0, r0, %f1388; mov.f32 %f1386, r0;} // end inline asm mov.u32 %r651, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1386, %r651, %r661, %r662; @p add.f32 r0, r0, %f1386; mov.f32 %f1389, r0;} // end inline asm mov.u32 %r654, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1389, %r654, %r661, %r662; @p add.f32 r0, r0, %f1389; mov.f32 %f1392, r0;} // end inline asm mov.u32 %r657, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1392, %r657, %r661, %r662; @p add.f32 r0, r0, %f1392; mov.f32 %f1395, r0;} // end inline asm mov.u32 %r660, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1395, %r660, %r661, %r662; @p add.f32 r0, r0, %f1395; mov.f32 %f1493, r0;} // end inline asm $L__BB0_29: @%p8 bra $L__BB0_31; add.s32 %r757, %r545, -112; st.shared.f32 [%r757+3200], %f1492; $L__BB0_31: bar.sync 0; @%p9 bra $L__BB0_33; ld.shared.f32 %f1421, [%r54+3216]; add.ftz.f32 %f1422, %f1492, %f1421; ld.shared.f32 %f1423, [%r54+3220]; add.ftz.f32 %f1424, %f1422, %f1423; ld.shared.f32 %f1425, [%r54+3224]; add.ftz.f32 %f1408, %f1424, %f1425; mov.u32 %r672, 1; mov.u32 %r685, 31; mov.u32 %r686, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1408, %r672, %r685, %r686; @p add.f32 r0, r0, %f1408; mov.f32 %f1406, r0;} // end inline asm mov.u32 %r675, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1406, %r675, %r685, %r686; @p add.f32 r0, r0, %f1406; mov.f32 %f1409, r0;} // end inline asm mov.u32 %r678, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1409, %r678, %r685, %r686; @p add.f32 r0, r0, %f1409; mov.f32 %f1412, r0;} // end inline asm mov.u32 %r681, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1412, %r681, %r685, %r686; @p add.f32 r0, r0, %f1412; mov.f32 %f1415, r0;} // end inline asm mov.u32 %r684, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1415, %r684, %r685, %r686; @p add.f32 r0, r0, %f1415; mov.f32 %f1492, r0;} // end inline asm $L__BB0_33: @%p8 bra $L__BB0_35; add.s32 %r758, %r545, -112; st.shared.f32 [%r758+3840], %f1491; $L__BB0_35: bar.sync 0; @%p9 bra $L__BB0_37; ld.shared.f32 %f1441, [%r54+3856]; add.ftz.f32 %f1442, %f1491, %f1441; ld.shared.f32 %f1443, [%r54+3860]; add.ftz.f32 %f1444, %f1442, %f1443; ld.shared.f32 %f1445, [%r54+3864]; add.ftz.f32 %f1428, %f1444, %f1445; mov.u32 %r696, 1; mov.u32 %r709, 31; mov.u32 %r710, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1428, %r696, %r709, %r710; @p add.f32 r0, r0, %f1428; mov.f32 %f1426, r0;} // end inline asm mov.u32 %r699, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1426, %r699, %r709, %r710; @p add.f32 r0, r0, %f1426; mov.f32 %f1429, r0;} // end inline asm mov.u32 %r702, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1429, %r702, %r709, %r710; @p add.f32 r0, r0, %f1429; mov.f32 %f1432, r0;} // end inline asm mov.u32 %r705, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1432, %r705, %r709, %r710; @p add.f32 r0, r0, %f1432; mov.f32 %f1435, r0;} // end inline asm mov.u32 %r708, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1435, %r708, %r709, %r710; @p add.f32 r0, r0, %f1435; mov.f32 %f1491, r0;} // end inline asm $L__BB0_37: @%p8 bra $L__BB0_39; add.s32 %r759, %r545, -112; st.shared.f32 [%r759+4480], %f1490; $L__BB0_39: bar.sync 0; @%p9 bra $L__BB0_41; ld.shared.f32 %f1461, [%r54+4496]; add.ftz.f32 %f1462, %f1490, %f1461; ld.shared.f32 %f1463, [%r54+4500]; add.ftz.f32 %f1464, %f1462, %f1463; ld.shared.f32 %f1465, [%r54+4504]; add.ftz.f32 %f1448, %f1464, %f1465; mov.u32 %r720, 1; mov.u32 %r733, 31; mov.u32 %r734, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1448, %r720, %r733, %r734; @p add.f32 r0, r0, %f1448; mov.f32 %f1446, r0;} // end inline asm mov.u32 %r723, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1446, %r723, %r733, %r734; @p add.f32 r0, r0, %f1446; mov.f32 %f1449, r0;} // end inline asm mov.u32 %r726, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1449, %r726, %r733, %r734; @p add.f32 r0, r0, %f1449; mov.f32 %f1452, r0;} // end inline asm mov.u32 %r729, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1452, %r729, %r733, %r734; @p add.f32 r0, r0, %f1452; mov.f32 %f1455, r0;} // end inline asm mov.u32 %r732, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f1455, %r732, %r733, %r734; @p add.f32 r0, r0, %f1455; mov.f32 %f1490, r0;} // end inline asm $L__BB0_41: or.b32 %r737, %r66, %r540; setp.ne.s32 %p24, %r737, 0; @%p24 bra $L__BB0_59; ld.param.u64 %rd107, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0+8]; setp.eq.s64 %p25, %rd107, 0; mul.ftz.f32 %f1514, %f95, %f1497; mov.u32 %r738, %ctaid.x; cvt.s64.s32 %rd15, %r738; @%p25 bra $L__BB0_44; shl.b64 %rd64, %rd15, 1; add.s64 %rd65, %rd2, %rd64; ld.global.u16 %rs763, [%rd65]; // begin inline asm { cvt.f32.f16 %f1466, %rs763;} // end inline asm fma.rn.ftz.f32 %f1514, %f96, %f1466, %f1514; $L__BB0_44: ld.param.u64 %rd108, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0]; // begin inline asm { cvt.rn.f16.f32 %rs764, %f1514;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd66, 1.0; // end inline asm shl.b64 %rd69, %rd15, 1; add.s64 %rd67, %rd108, %rd69; // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd67], %rs764, %rd66; // end inline asm mul.ftz.f32 %f1515, %f95, %f1496; add.s32 %r740, %r62, %r738; cvt.s64.s32 %rd18, %r740; @%p25 bra $L__BB0_46; shl.b64 %rd70, %rd18, 1; add.s64 %rd71, %rd2, %rd70; ld.global.u16 %rs766, [%rd71]; // begin inline asm { cvt.f32.f16 %f1468, %rs766;} // end inline asm fma.rn.ftz.f32 %f1515, %f96, %f1468, %f1515; $L__BB0_46: cvt.s64.s32 %rd19, %r62; mul.wide.s32 %rd75, %r62, 2; add.s64 %rd73, %rd67, %rd75; // begin inline asm { cvt.rn.f16.f32 %rs767, %f1515;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd72, 1.0; // end inline asm // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd73], %rs767, %rd72; // end inline asm mul.ftz.f32 %f1516, %f95, %f1495; cvt.u32.u64 %r741, %rd18; add.s32 %r742, %r741, %r62; cvt.s64.s32 %rd20, %r742; @%p25 bra $L__BB0_48; shl.b64 %rd76, %rd20, 1; add.s64 %rd77, %rd2, %rd76; ld.global.u16 %rs769, [%rd77]; // begin inline asm { cvt.f32.f16 %f1470, %rs769;} // end inline asm fma.rn.ftz.f32 %f1516, %f96, %f1470, %f1516; $L__BB0_48: ld.param.u64 %rd109, [_Z28dequant_gemv_group128_batch823DequantGemvKernelParams_param_0]; shl.b64 %rd81, %rd20, 1; add.s64 %rd79, %rd109, %rd81; // begin inline asm { cvt.rn.f16.f32 %rs770, %f1516;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd78, 1.0; // end inline asm // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd79], %rs770, %rd78; // end inline asm mul.ftz.f32 %f1517, %f95, %f1494; cvt.u32.u64 %r743, %rd20; add.s32 %r744, %r743, %r62; cvt.s64.s32 %rd22, %r744; @%p25 bra $L__BB0_50; shl.b64 %rd82, %rd22, 1; add.s64 %rd83, %rd2, %rd82; ld.global.u16 %rs772, [%rd83]; // begin inline asm { cvt.f32.f16 %f1472, %rs772;} // end inline asm fma.rn.ftz.f32 %f1517, %f96, %f1472, %f1517; $L__BB0_50: // begin inline asm { cvt.rn.f16.f32 %rs773, %f1517;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd84, 1.0; // end inline asm shl.b64 %rd23, %rd19, 1; add.s64 %rd85, %rd79, %rd23; // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd85], %rs773, %rd84; // end inline asm mul.ftz.f32 %f1518, %f95, %f1493; cvt.u32.u64 %r745, %rd22; add.s32 %r746, %r745, %r62; cvt.s64.s32 %rd25, %r746; @%p25 bra $L__BB0_52; shl.b64 %rd87, %rd25, 1; add.s64 %rd88, %rd2, %rd87; ld.global.u16 %rs775, [%rd88]; // begin inline asm { cvt.f32.f16 %f1474, %rs775;} // end inline asm fma.rn.ftz.f32 %f1518, %f96, %f1474, %f1518; $L__BB0_52: // begin inline asm { cvt.rn.f16.f32 %rs776, %f1518;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd89, 1.0; // end inline asm add.s64 %rd90, %rd85, %rd23; // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd90], %rs776, %rd89; // end inline asm mul.ftz.f32 %f1519, %f95, %f1492; cvt.u32.u64 %r747, %rd25; add.s32 %r748, %r747, %r62; cvt.s64.s32 %rd27, %r748; @%p25 bra $L__BB0_54; shl.b64 %rd92, %rd27, 1; add.s64 %rd93, %rd2, %rd92; ld.global.u16 %rs778, [%rd93]; // begin inline asm { cvt.f32.f16 %f1476, %rs778;} // end inline asm fma.rn.ftz.f32 %f1519, %f96, %f1476, %f1519; $L__BB0_54: // begin inline asm { cvt.rn.f16.f32 %rs779, %f1519;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd94, 1.0; // end inline asm add.s64 %rd95, %rd90, %rd23; // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd95], %rs779, %rd94; // end inline asm mul.ftz.f32 %f1520, %f95, %f1491; cvt.u32.u64 %r749, %rd27; add.s32 %r750, %r749, %r62; cvt.s64.s32 %rd29, %r750; @%p25 bra $L__BB0_56; shl.b64 %rd97, %rd29, 1; add.s64 %rd98, %rd2, %rd97; ld.global.u16 %rs781, [%rd98]; // begin inline asm { cvt.f32.f16 %f1478, %rs781;} // end inline asm fma.rn.ftz.f32 %f1520, %f96, %f1478, %f1520; $L__BB0_56: // begin inline asm { cvt.rn.f16.f32 %rs782, %f1520;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd99, 1.0; // end inline asm add.s64 %rd100, %rd95, %rd23; // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd100], %rs782, %rd99; // end inline asm mul.ftz.f32 %f1521, %f95, %f1490; cvt.u32.u64 %r751, %rd29; add.s32 %r56, %r751, %r62; @%p25 bra $L__BB0_58; mul.wide.s32 %rd102, %r56, 2; add.s64 %rd103, %rd2, %rd102; ld.global.u16 %rs784, [%rd103]; // begin inline asm { cvt.f32.f16 %f1480, %rs784;} // end inline asm fma.rn.ftz.f32 %f1521, %f96, %f1480, %f1521; $L__BB0_58: // begin inline asm { cvt.rn.f16.f32 %rs785, %f1521;} // end inline asm // begin inline asm createpolicy.fractional.L2::evict_last.L2::evict_unchanged.b64 %rd104, 1.0; // end inline asm add.s64 %rd105, %rd100, %rd23; // begin inline asm st.global.L1::evict_normal.L2::cache_hint.u16 [%rd105], %rs785, %rd104; // end inline asm $L__BB0_59: ret; } // .globl _ZN3cub11EmptyKernelIvEEvv .visible .entry _ZN3cub11EmptyKernelIvEEvv() { ret; }