global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders2_4E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders2_5E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders2_6E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders2_7E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders2_8E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders2_9E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_2882efee6thrust12placeholders3_10E[1]; .visible .entry _Z27dequant_gemv_group64_batch723DequantGemvKernelParams( .param .align 8 .b8 _Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0[80] ) { .reg .pred %p<24>; .reg .b16 %rs<394>; .reg .f32 %f<797>; .reg .b32 %r<371>; .reg .b64 %rd<69>; // demoted variable .shared .align 16 .b8 _ZZ9gemv_int4ILi4ELi64ELi7EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage[4480]; ld.param.v2.u32 {%r52, %r53}, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+48]; ld.param.v2.u32 {%r54, %r55}, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+56]; ld.param.v2.f32 {%f68, %f69}, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+64]; ld.param.v4.u8 {%rs37, %rs38, %rs39, %rs40}, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+72]; ld.param.u64 %rd25, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+40]; ld.param.u64 %rd24, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+32]; ld.param.u64 %rd23, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+24]; ld.param.u64 %rd22, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+16]; mov.u32 %r370, %tid.y; shl.b32 %r56, %r370, 5; mov.u32 %r57, %tid.x; add.s32 %r369, %r56, %r57; shl.b32 %r368, %r369, 1; setp.ge.u32 %p1, %r368, %r54; mov.f32 %f776, 0f00000000; mov.f32 %f777, %f776; mov.f32 %f778, %f776; mov.f32 %f779, %f776; mov.f32 %f780, %f776; mov.f32 %f781, %f776; mov.f32 %f782, %f776; @%p1 bra $L__BB0_9; cvta.to.global.u64 %rd2, %rd22; mov.u32 %r58, %ctaid.x; mul.lo.s32 %r6, %r55, %r58; shl.b16 %rs2, %rs37, 3; cvta.to.global.u64 %rd3, %rd23; cvta.to.global.u64 %rd4, %rd25; $L__BB0_2: mad.lo.s32 %r60, %r54, %r58, %r368; mul.wide.u32 %rd26, %r60, 4; add.s64 %rd27, %rd3, %rd26; ld.global.v2.u32 {%r61, %r62}, [%rd27]; shr.u32 %r64, %r57, 2; shl.b32 %r65, %r370, 3; add.s32 %r12, %r65, %r64; add.s32 %r13, %r12, %r6; mul.wide.s32 %rd28, %r13, 2; add.s64 %rd29, %rd4, %rd28; ld.global.u16 %rs45, [%rd29]; // begin inline asm { cvt.f32.f16 %f77, %rs45;} // end inline asm setp.eq.s64 %p2, %rd24, 0; mov.u16 %rs393, %rs2; @%p2 bra $L__BB0_4; shr.u32 %r66, %r13, 31; add.s32 %r67, %r13, %r66; shr.s32 %r68, %r67, 1; cvt.s64.s32 %rd30, %r68; cvta.to.global.u64 %rd31, %rd24; add.s64 %rd32, %rd31, %rd30; ld.global.u8 %r69, [%rd32]; shl.b32 %r70, %r12, 2; and.b32 %r71, %r70, 4; shr.u32 %r72, %r69, %r71; cvt.u16.u32 %rs46, %r72; and.b16 %rs393, %rs46, 15; $L__BB0_4: shl.b32 %r14, %r369, 4; setp.ge.s32 %p3, %r14, %r52; @%p3 bra $L__BB0_8; setp.eq.s16 %p4, %rs37, 0; shr.u16 %rs48, %rs393, 3; and.b16 %rs49, %rs48, 1; setp.eq.b16 %p5, %rs49, 1; and.pred %p6, %p4, %p5; selp.b16 %rs50, -16, 0, %p6; or.b16 %rs51, %rs50, %rs393; cvt.s16.s8 %rs52, %rs51; cvt.rn.f32.s16 %f9, %rs52; mul.wide.s32 %rd33, %r14, 2; add.s64 %rd5, %rd2, %rd33; ld.global.v4.u32 {%r73, %r74, %r75, %r76}, [%rd5]; mul.wide.s32 %rd34, %r52, 2; add.s64 %rd6, %rd5, %rd34; ld.global.v4.u32 {%r77, %r78, %r79, %r80}, [%rd6]; add.s32 %r81, %r14, %r52; add.s32 %r82, %r81, %r52; mul.wide.s32 %rd35, %r82, 2; add.s64 %rd36, %rd2, %rd35; ld.global.v4.u32 {%r83, %r84, %r85, %r86}, [%rd36]; add.s64 %rd37, %rd36, %rd34; ld.global.v4.u32 {%r87, %r88, %r89, %r90}, [%rd37]; add.s64 %rd38, %rd37, %rd34; ld.global.v4.u32 {%r91, %r92, %r93, %r94}, [%rd38]; add.s64 %rd39, %rd38, %rd34; ld.global.v4.u32 {%r95, %r96, %r97, %r98}, [%rd39]; add.s64 %rd40, %rd39, %rd34; ld.global.v4.u32 {%r99, %r100, %r101, %r102}, [%rd40]; cvt.u16.u32 %rs53, %r61; shr.u16 %rs54, %rs53, 4; and.b16 %rs5, %rs54, 15; mov.b32 {%rs55, %rs6}, %r73; mov.b32 {%rs56, %rs7}, %r77; mov.b32 {%rs57, %rs8}, %r83; shr.u16 %rs58, %rs53, 8; and.b16 %rs9, %rs58, 15; shr.u16 %rs10, %rs53, 12; shr.u32 %r103, %r61, 16; cvt.u16.u32 %rs59, %r103; and.b16 %rs11, %rs59, 15; shr.u32 %r104, %r61, 20; cvt.u16.u32 %rs60, %r104; and.b16 %rs12, %rs60, 15; shr.u32 %r105, %r61, 24; cvt.u16.u32 %rs61, %r105; and.b16 %rs13, %rs61, 15; shr.u32 %r106, %r61, 28; cvt.u16.u32 %rs14, %r106; cvt.u16.u32 %rs62, %r62; and.b16 %rs15, %rs62, 15; shr.u32 %r107, %r62, 4; cvt.u16.u32 %rs16, %r107; and.b16 %rs17, %rs16, 15; shr.u32 %r108, %r62, 8; cvt.u16.u32 %rs18, %r108; and.b16 %rs19, %rs18, 15; shr.u32 %r109, %r62, 12; cvt.u16.u32 %rs20, %r109; and.b16 %rs21, %rs20, 15; shr.u32 %r110, %r62, 16; cvt.u16.u32 %rs22, %r110; and.b16 %rs23, %rs22, 15; shr.u32 %r111, %r62, 20; cvt.u16.u32 %rs24, %r111; and.b16 %rs25, %rs24, 15; shr.u32 %r112, %r62, 24; cvt.u16.u32 %rs26, %r112; and.b16 %rs27, %rs26, 15; shr.u32 %r113, %r62, 28; cvt.u16.u32 %rs28, %r113; add.s64 %rd41, %rd6, %rd34; add.s64 %rd7, %rd41, 16; add.s64 %rd8, %rd7, %rd34; add.s64 %rd9, %rd8, %rd34; add.s64 %rd10, %rd9, %rd34; add.s64 %rd11, %rd10, %rd34; @%p4 bra $L__BB0_7; and.b16 %rs176, %rs53, 15; cvt.rn.f32.s16 %f190, %rs176; sub.ftz.f32 %f191, %f190, %f9; mul.ftz.f32 %f192, %f77, %f191; cvt.u16.u32 %rs63, %r73; // begin inline asm { cvt.f32.f16 %f78, %rs63;} // end inline asm fma.rn.ftz.f32 %f193, %f192, %f78, %f782; cvt.u16.u32 %rs64, %r77; // begin inline asm { cvt.f32.f16 %f79, %rs64;} // end inline asm fma.rn.ftz.f32 %f194, %f192, %f79, %f781; cvt.u16.u32 %rs65, %r83; // begin inline asm { cvt.f32.f16 %f80, %rs65;} // end inline asm fma.rn.ftz.f32 %f195, %f192, %f80, %f780; mov.b32 {%rs66, %rs73}, %r87; // begin inline asm { cvt.f32.f16 %f81, %rs66;} // end inline asm fma.rn.ftz.f32 %f196, %f192, %f81, %f779; mov.b32 {%rs67, %rs74}, %r91; // begin inline asm { cvt.f32.f16 %f82, %rs67;} // end inline asm fma.rn.ftz.f32 %f197, %f192, %f82, %f778; mov.b32 {%rs68, %rs75}, %r95; // begin inline asm { cvt.f32.f16 %f83, %rs68;} // end inline asm fma.rn.ftz.f32 %f198, %f192, %f83, %f777; mov.b32 {%rs69, %rs76}, %r99; // begin inline asm { cvt.f32.f16 %f84, %rs69;} // end inline asm fma.rn.ftz.f32 %f199, %f192, %f84, %f776; cvt.rn.f32.s16 %f200, %rs5; sub.ftz.f32 %f201, %f200, %f9; mul.ftz.f32 %f202, %f77, %f201; // begin inline asm { cvt.f32.f16 %f85, %rs6;} // end inline asm fma.rn.ftz.f32 %f203, %f202, %f85, %f193; // begin inline asm { cvt.f32.f16 %f86, %rs7;} // end inline asm fma.rn.ftz.f32 %f204, %f202, %f86, %f194; // begin inline asm { cvt.f32.f16 %f87, %rs8;} // end inline asm fma.rn.ftz.f32 %f205, %f202, %f87, %f195; // begin inline asm { cvt.f32.f16 %f88, %rs73;} // end inline asm fma.rn.ftz.f32 %f206, %f202, %f88, %f196; // begin inline asm { cvt.f32.f16 %f89, %rs74;} // end inline asm fma.rn.ftz.f32 %f207, %f202, %f89, %f197; // begin inline asm { cvt.f32.f16 %f90, %rs75;} // end inline asm fma.rn.ftz.f32 %f208, %f202, %f90, %f198; // begin inline asm { cvt.f32.f16 %f91, %rs76;} // end inline asm fma.rn.ftz.f32 %f209, %f202, %f91, %f199; cvt.rn.f32.s16 %f210, %rs9; sub.ftz.f32 %f211, %f210, %f9; mul.ftz.f32 %f212, %f77, %f211; mov.b32 {%rs77, %rs84}, %r74; // begin inline asm { cvt.f32.f16 %f92, %rs77;} // end inline asm fma.rn.ftz.f32 %f213, %f212, %f92, %f203; mov.b32 {%rs78, %rs85}, %r78; // begin inline asm { cvt.f32.f16 %f93, %rs78;} // end inline asm fma.rn.ftz.f32 %f214, %f212, %f93, %f204; mov.b32 {%rs79, %rs86}, %r84; // begin inline asm { cvt.f32.f16 %f94, %rs79;} // end inline asm fma.rn.ftz.f32 %f215, %f212, %f94, %f205; mov.b32 {%rs80, %rs87}, %r88; // begin inline asm { cvt.f32.f16 %f95, %rs80;} // end inline asm fma.rn.ftz.f32 %f216, %f212, %f95, %f206; mov.b32 {%rs81, %rs88}, %r92; // begin inline asm { cvt.f32.f16 %f96, %rs81;} // end inline asm fma.rn.ftz.f32 %f217, %f212, %f96, %f207; mov.b32 {%rs82, %rs89}, %r96; // begin inline asm { cvt.f32.f16 %f97, %rs82;} // end inline asm fma.rn.ftz.f32 %f218, %f212, %f97, %f208; mov.b32 {%rs83, %rs90}, %r100; // begin inline asm { cvt.f32.f16 %f98, %rs83;} // end inline asm fma.rn.ftz.f32 %f219, %f212, %f98, %f209; cvt.rn.f32.s16 %f220, %rs10; sub.ftz.f32 %f221, %f220, %f9; mul.ftz.f32 %f222, %f77, %f221; // begin inline asm { cvt.f32.f16 %f99, %rs84;} // end inline asm fma.rn.ftz.f32 %f223, %f222, %f99, %f213; // begin inline asm { cvt.f32.f16 %f100, %rs85;} // end inline asm fma.rn.ftz.f32 %f224, %f222, %f100, %f214; // begin inline asm { cvt.f32.f16 %f101, %rs86;} // end inline asm fma.rn.ftz.f32 %f225, %f222, %f101, %f215; // begin inline asm { cvt.f32.f16 %f102, %rs87;} // end inline asm fma.rn.ftz.f32 %f226, %f222, %f102, %f216; // begin inline asm { cvt.f32.f16 %f103, %rs88;} // end inline asm fma.rn.ftz.f32 %f227, %f222, %f103, %f217; // begin inline asm { cvt.f32.f16 %f104, %rs89;} // end inline asm fma.rn.ftz.f32 %f228, %f222, %f104, %f218; // begin inline asm { cvt.f32.f16 %f105, %rs90;} // end inline asm fma.rn.ftz.f32 %f229, %f222, %f105, %f219; cvt.rn.f32.s16 %f230, %rs11; sub.ftz.f32 %f231, %f230, %f9; mul.ftz.f32 %f232, %f77, %f231; mov.b32 {%rs91, %rs98}, %r75; // begin inline asm { cvt.f32.f16 %f106, %rs91;} // end inline asm fma.rn.ftz.f32 %f233, %f232, %f106, %f223; mov.b32 {%rs92, %rs99}, %r79; // begin inline asm { cvt.f32.f16 %f107, %rs92;} // end inline asm fma.rn.ftz.f32 %f234, %f232, %f107, %f224; mov.b32 {%rs93, %rs100}, %r85; // begin inline asm { cvt.f32.f16 %f108, %rs93;} // end inline asm fma.rn.ftz.f32 %f235, %f232, %f108, %f225; mov.b32 {%rs94, %rs101}, %r89; // begin inline asm { cvt.f32.f16 %f109, %rs94;} // end inline asm fma.rn.ftz.f32 %f236, %f232, %f109, %f226; mov.b32 {%rs95, %rs102}, %r93; // begin inline asm { cvt.f32.f16 %f110, %rs95;} // end inline asm fma.rn.ftz.f32 %f237, %f232, %f110, %f227; mov.b32 {%rs96, %rs103}, %r97; // begin inline asm { cvt.f32.f16 %f111, %rs96;} // end inline asm fma.rn.ftz.f32 %f238, %f232, %f111, %f228; mov.b32 {%rs97, %rs104}, %r101; // begin inline asm { cvt.f32.f16 %f112, %rs97;} // end inline asm fma.rn.ftz.f32 %f239, %f232, %f112, %f229; cvt.rn.f32.s16 %f240, %rs12; sub.ftz.f32 %f241, %f240, %f9; mul.ftz.f32 %f242, %f77, %f241; // begin inline asm { cvt.f32.f16 %f113, %rs98;} // end inline asm fma.rn.ftz.f32 %f243, %f242, %f113, %f233; // begin inline asm { cvt.f32.f16 %f114, %rs99;} // end inline asm fma.rn.ftz.f32 %f244, %f242, %f114, %f234; // begin inline asm { cvt.f32.f16 %f115, %rs100;} // end inline asm fma.rn.ftz.f32 %f245, %f242, %f115, %f235; // begin inline asm { cvt.f32.f16 %f116, %rs101;} // end inline asm fma.rn.ftz.f32 %f246, %f242, %f116, %f236; // begin inline asm { cvt.f32.f16 %f117, %rs102;} // end inline asm fma.rn.ftz.f32 %f247, %f242, %f117, %f237; // begin inline asm { cvt.f32.f16 %f118, %rs103;} // end inline asm fma.rn.ftz.f32 %f248, %f242, %f118, %f238; // begin inline asm { cvt.f32.f16 %f119, %rs104;} // end inline asm fma.rn.ftz.f32 %f249, %f242, %f119, %f239; cvt.rn.f32.s16 %f250, %rs13; sub.ftz.f32 %f251, %f250, %f9; mul.ftz.f32 %f252, %f77, %f251; mov.b32 {%rs105, %rs112}, %r76; // begin inline asm { cvt.f32.f16 %f120, %rs105;} // end inline asm fma.rn.ftz.f32 %f253, %f252, %f120, %f243; mov.b32 {%rs106, %rs113}, %r80; // begin inline asm { cvt.f32.f16 %f121, %rs106;} // end inline asm fma.rn.ftz.f32 %f254, %f252, %f121, %f244; mov.b32 {%rs107, %rs114}, %r86; // begin inline asm { cvt.f32.f16 %f122, %rs107;} // end inline asm fma.rn.ftz.f32 %f255, %f252, %f122, %f245; mov.b32 {%rs108, %rs115}, %r90; // begin inline asm { cvt.f32.f16 %f123, %rs108;} // end inline asm fma.rn.ftz.f32 %f256, %f252, %f123, %f246; mov.b32 {%rs109, %rs116}, %r94; // begin inline asm { cvt.f32.f16 %f124, %rs109;} // end inline asm fma.rn.ftz.f32 %f257, %f252, %f124, %f247; mov.b32 {%rs110, %rs117}, %r98; // begin inline asm { cvt.f32.f16 %f125, %rs110;} // end inline asm fma.rn.ftz.f32 %f258, %f252, %f125, %f248; mov.b32 {%rs111, %rs118}, %r102; // begin inline asm { cvt.f32.f16 %f126, %rs111;} // end inline asm fma.rn.ftz.f32 %f259, %f252, %f126, %f249; cvt.rn.f32.s16 %f260, %rs14; sub.ftz.f32 %f261, %f260, %f9; mul.ftz.f32 %f262, %f77, %f261; // begin inline asm { cvt.f32.f16 %f127, %rs112;} // end inline asm fma.rn.ftz.f32 %f263, %f262, %f127, %f253; // begin inline asm { cvt.f32.f16 %f128, %rs113;} // end inline asm fma.rn.ftz.f32 %f264, %f262, %f128, %f254; // begin inline asm { cvt.f32.f16 %f129, %rs114;} // end inline asm fma.rn.ftz.f32 %f265, %f262, %f129, %f255; // begin inline asm { cvt.f32.f16 %f130, %rs115;} // end inline asm fma.rn.ftz.f32 %f266, %f262, %f130, %f256; // begin inline asm { cvt.f32.f16 %f131, %rs116;} // end inline asm fma.rn.ftz.f32 %f267, %f262, %f131, %f257; // begin inline asm { cvt.f32.f16 %f132, %rs117;} // end inline asm fma.rn.ftz.f32 %f268, %f262, %f132, %f258; // begin inline asm { cvt.f32.f16 %f133, %rs118;} // end inline asm fma.rn.ftz.f32 %f269, %f262, %f133, %f259; ld.global.v4.u32 {%r114, %r115, %r116, %r117}, [%rd5+16]; ld.global.v4.u32 {%r122, %r123, %r124, %r125}, [%rd6+16]; ld.global.v4.u32 {%r130, %r131, %r132, %r133}, [%rd7]; ld.global.v4.u32 {%r138, %r139, %r140, %r141}, [%rd8]; ld.global.v4.u32 {%r146, %r147, %r148, %r149}, [%rd9]; ld.global.v4.u32 {%r154, %r155, %r156, %r157}, [%rd10]; ld.global.v4.u32 {%r162, %r163, %r164, %r165}, [%rd11]; cvt.rn.f32.s16 %f270, %rs15; sub.ftz.f32 %f271, %f270, %f9; mul.ftz.f32 %f272, %f77, %f271; mov.b32 {%rs119, %rs126}, %r114; // begin inline asm { cvt.f32.f16 %f134, %rs119;} // end inline asm fma.rn.ftz.f32 %f273, %f272, %f134, %f263; mov.b32 {%rs120, %rs127}, %r122; // begin inline asm { cvt.f32.f16 %f135, %rs120;} // end inline asm fma.rn.ftz.f32 %f274, %f272, %f135, %f264; mov.b32 {%rs121, %rs128}, %r130; // begin inline asm { cvt.f32.f16 %f136, %rs121;} // end inline asm fma.rn.ftz.f32 %f275, %f272, %f136, %f265; mov.b32 {%rs122, %rs129}, %r138; // begin inline asm { cvt.f32.f16 %f137, %rs122;} // end inline asm fma.rn.ftz.f32 %f276, %f272, %f137, %f266; mov.b32 {%rs123, %rs130}, %r146; // begin inline asm { cvt.f32.f16 %f138, %rs123;} // end inline asm fma.rn.ftz.f32 %f277, %f272, %f138, %f267; mov.b32 {%rs124, %rs131}, %r154; // begin inline asm { cvt.f32.f16 %f139, %rs124;} // end inline asm fma.rn.ftz.f32 %f278, %f272, %f139, %f268; mov.b32 {%rs125, %rs132}, %r162; // begin inline asm { cvt.f32.f16 %f140, %rs125;} // end inline asm fma.rn.ftz.f32 %f279, %f272, %f140, %f269; cvt.rn.f32.s16 %f280, %rs17; sub.ftz.f32 %f281, %f280, %f9; mul.ftz.f32 %f282, %f77, %f281; // begin inline asm { cvt.f32.f16 %f141, %rs126;} // end inline asm fma.rn.ftz.f32 %f283, %f282, %f141, %f273; // begin inline asm { cvt.f32.f16 %f142, %rs127;} // end inline asm fma.rn.ftz.f32 %f284, %f282, %f142, %f274; // begin inline asm { cvt.f32.f16 %f143, %rs128;} // end inline asm fma.rn.ftz.f32 %f285, %f282, %f143, %f275; // begin inline asm { cvt.f32.f16 %f144, %rs129;} // end inline asm fma.rn.ftz.f32 %f286, %f282, %f144, %f276; // begin inline asm { cvt.f32.f16 %f145, %rs130;} // end inline asm fma.rn.ftz.f32 %f287, %f282, %f145, %f277; // begin inline asm { cvt.f32.f16 %f146, %rs131;} // end inline asm fma.rn.ftz.f32 %f288, %f282, %f146, %f278; // begin inline asm { cvt.f32.f16 %f147, %rs132;} // end inline asm fma.rn.ftz.f32 %f289, %f282, %f147, %f279; cvt.rn.f32.s16 %f290, %rs19; sub.ftz.f32 %f291, %f290, %f9; mul.ftz.f32 %f292, %f77, %f291; mov.b32 {%rs133, %rs140}, %r115; // begin inline asm { cvt.f32.f16 %f148, %rs133;} // end inline asm fma.rn.ftz.f32 %f293, %f292, %f148, %f283; mov.b32 {%rs134, %rs141}, %r123; // begin inline asm { cvt.f32.f16 %f149, %rs134;} // end inline asm fma.rn.ftz.f32 %f294, %f292, %f149, %f284; mov.b32 {%rs135, %rs142}, %r131; // begin inline asm { cvt.f32.f16 %f150, %rs135;} // end inline asm fma.rn.ftz.f32 %f295, %f292, %f150, %f285; mov.b32 {%rs136, %rs143}, %r139; // begin inline asm { cvt.f32.f16 %f151, %rs136;} // end inline asm fma.rn.ftz.f32 %f296, %f292, %f151, %f286; mov.b32 {%rs137, %rs144}, %r147; // begin inline asm { cvt.f32.f16 %f152, %rs137;} // end inline asm fma.rn.ftz.f32 %f297, %f292, %f152, %f287; mov.b32 {%rs138, %rs145}, %r155; // begin inline asm { cvt.f32.f16 %f153, %rs138;} // end inline asm fma.rn.ftz.f32 %f298, %f292, %f153, %f288; mov.b32 {%rs139, %rs146}, %r163; // begin inline asm { cvt.f32.f16 %f154, %rs139;} // end inline asm fma.rn.ftz.f32 %f299, %f292, %f154, %f289; cvt.rn.f32.s16 %f300, %rs21; sub.ftz.f32 %f301, %f300, %f9; mul.ftz.f32 %f302, %f77, %f301; // begin inline asm { cvt.f32.f16 %f155, %rs140;} // end inline asm fma.rn.ftz.f32 %f303, %f302, %f155, %f293; // begin inline asm { cvt.f32.f16 %f156, %rs141;} // end inline asm fma.rn.ftz.f32 %f304, %f302, %f156, %f294; // begin inline asm { cvt.f32.f16 %f157, %rs142;} // end inline asm fma.rn.ftz.f32 %f305, %f302, %f157, %f295; // begin inline asm { cvt.f32.f16 %f158, %rs143;} // end inline asm fma.rn.ftz.f32 %f306, %f302, %f158, %f296; // begin inline asm { cvt.f32.f16 %f159, %rs144;} // end inline asm fma.rn.ftz.f32 %f307, %f302, %f159, %f297; // begin inline asm { cvt.f32.f16 %f160, %rs145;} // end inline asm fma.rn.ftz.f32 %f308, %f302, %f160, %f298; // begin inline asm { cvt.f32.f16 %f161, %rs146;} // end inline asm fma.rn.ftz.f32 %f309, %f302, %f161, %f299; cvt.rn.f32.s16 %f310, %rs23; sub.ftz.f32 %f311, %f310, %f9; mul.ftz.f32 %f312, %f77, %f311; mov.b32 {%rs147, %rs154}, %r116; // begin inline asm { cvt.f32.f16 %f162, %rs147;} // end inline asm fma.rn.ftz.f32 %f313, %f312, %f162, %f303; mov.b32 {%rs148, %rs155}, %r124; // begin inline asm { cvt.f32.f16 %f163, %rs148;} // end inline asm fma.rn.ftz.f32 %f314, %f312, %f163, %f304; mov.b32 {%rs149, %rs156}, %r132; // begin inline asm { cvt.f32.f16 %f164, %rs149;} // end inline asm fma.rn.ftz.f32 %f315, %f312, %f164, %f305; mov.b32 {%rs150, %rs157}, %r140; // begin inline asm { cvt.f32.f16 %f165, %rs150;} // end inline asm fma.rn.ftz.f32 %f316, %f312, %f165, %f306; mov.b32 {%rs151, %rs158}, %r148; // begin inline asm { cvt.f32.f16 %f166, %rs151;} // end inline asm fma.rn.ftz.f32 %f317, %f312, %f166, %f307; mov.b32 {%rs152, %rs159}, %r156; // begin inline asm { cvt.f32.f16 %f167, %rs152;} // end inline asm fma.rn.ftz.f32 %f318, %f312, %f167, %f308; mov.b32 {%rs153, %rs160}, %r164; // begin inline asm { cvt.f32.f16 %f168, %rs153;} // end inline asm fma.rn.ftz.f32 %f319, %f312, %f168, %f309; cvt.rn.f32.s16 %f320, %rs25; sub.ftz.f32 %f321, %f320, %f9; mul.ftz.f32 %f322, %f77, %f321; // begin inline asm { cvt.f32.f16 %f169, %rs154;} // end inline asm fma.rn.ftz.f32 %f323, %f322, %f169, %f313; // begin inline asm { cvt.f32.f16 %f170, %rs155;} // end inline asm fma.rn.ftz.f32 %f324, %f322, %f170, %f314; // begin inline asm { cvt.f32.f16 %f171, %rs156;} // end inline asm fma.rn.ftz.f32 %f325, %f322, %f171, %f315; // begin inline asm { cvt.f32.f16 %f172, %rs157;} // end inline asm fma.rn.ftz.f32 %f326, %f322, %f172, %f316; // begin inline asm { cvt.f32.f16 %f173, %rs158;} // end inline asm fma.rn.ftz.f32 %f327, %f322, %f173, %f317; // begin inline asm { cvt.f32.f16 %f174, %rs159;} // end inline asm fma.rn.ftz.f32 %f328, %f322, %f174, %f318; // begin inline asm { cvt.f32.f16 %f175, %rs160;} // end inline asm fma.rn.ftz.f32 %f329, %f322, %f175, %f319; cvt.rn.f32.s16 %f330, %rs27; sub.ftz.f32 %f331, %f330, %f9; mul.ftz.f32 %f332, %f77, %f331; mov.b32 {%rs161, %rs168}, %r117; // begin inline asm { cvt.f32.f16 %f176, %rs161;} // end inline asm fma.rn.ftz.f32 %f333, %f332, %f176, %f323; mov.b32 {%rs162, %rs169}, %r125; // begin inline asm { cvt.f32.f16 %f177, %rs162;} // end inline asm fma.rn.ftz.f32 %f334, %f332, %f177, %f324; mov.b32 {%rs163, %rs170}, %r133; // begin inline asm { cvt.f32.f16 %f178, %rs163;} // end inline asm fma.rn.ftz.f32 %f335, %f332, %f178, %f325; mov.b32 {%rs164, %rs171}, %r141; // begin inline asm { cvt.f32.f16 %f179, %rs164;} // end inline asm fma.rn.ftz.f32 %f336, %f332, %f179, %f326; mov.b32 {%rs165, %rs172}, %r149; // begin inline asm { cvt.f32.f16 %f180, %rs165;} // end inline asm fma.rn.ftz.f32 %f337, %f332, %f180, %f327; mov.b32 {%rs166, %rs173}, %r157; // begin inline asm { cvt.f32.f16 %f181, %rs166;} // end inline asm fma.rn.ftz.f32 %f338, %f332, %f181, %f328; mov.b32 {%rs167, %rs174}, %r165; // begin inline asm { cvt.f32.f16 %f182, %rs167;} // end inline asm fma.rn.ftz.f32 %f339, %f332, %f182, %f329; cvt.rn.f32.s16 %f340, %rs28; sub.ftz.f32 %f341, %f340, %f9; mul.ftz.f32 %f342, %f77, %f341; // begin inline asm { cvt.f32.f16 %f183, %rs168;} // end inline asm fma.rn.ftz.f32 %f782, %f342, %f183, %f333; // begin inline asm { cvt.f32.f16 %f184, %rs169;} // end inline asm fma.rn.ftz.f32 %f781, %f342, %f184, %f334; // begin inline asm { cvt.f32.f16 %f185, %rs170;} // end inline asm fma.rn.ftz.f32 %f780, %f342, %f185, %f335; // begin inline asm { cvt.f32.f16 %f186, %rs171;} // end inline asm fma.rn.ftz.f32 %f779, %f342, %f186, %f336; // begin inline asm { cvt.f32.f16 %f187, %rs172;} // end inline asm fma.rn.ftz.f32 %f778, %f342, %f187, %f337; // begin inline asm { cvt.f32.f16 %f188, %rs173;} // end inline asm fma.rn.ftz.f32 %f777, %f342, %f188, %f338; // begin inline asm { cvt.f32.f16 %f189, %rs174;} // end inline asm fma.rn.ftz.f32 %f776, %f342, %f189, %f339; bra.uni $L__BB0_8; $L__BB0_7: shl.b16 %rs290, %rs53, 4; cvt.s16.s8 %rs291, %rs290; shr.s16 %rs292, %rs291, 7; and.b16 %rs293, %rs292, -16; and.b16 %rs294, %rs53, 15; or.b16 %rs295, %rs293, %rs294; cvt.rn.f32.s16 %f455, %rs295; sub.ftz.f32 %f456, %f455, %f9; mul.ftz.f32 %f457, %f77, %f456; cvt.u16.u32 %rs177, %r73; // begin inline asm { cvt.f32.f16 %f343, %rs177;} // end inline asm fma.rn.ftz.f32 %f458, %f457, %f343, %f782; cvt.u16.u32 %rs178, %r77; // begin inline asm { cvt.f32.f16 %f344, %rs178;} // end inline asm fma.rn.ftz.f32 %f459, %f457, %f344, %f781; cvt.u16.u32 %rs179, %r83; // begin inline asm { cvt.f32.f16 %f345, %rs179;} // end inline asm fma.rn.ftz.f32 %f460, %f457, %f345, %f780; mov.b32 {%rs180, %rs187}, %r87; // begin inline asm { cvt.f32.f16 %f346, %rs180;} // end inline asm fma.rn.ftz.f32 %f461, %f457, %f346, %f779; mov.b32 {%rs181, %rs188}, %r91; // begin inline asm { cvt.f32.f16 %f347, %rs181;} // end inline asm fma.rn.ftz.f32 %f462, %f457, %f347, %f778; mov.b32 {%rs182, %rs189}, %r95; // begin inline asm { cvt.f32.f16 %f348, %rs182;} // end inline asm fma.rn.ftz.f32 %f463, %f457, %f348, %f777; mov.b32 {%rs183, %rs190}, %r99; // begin inline asm { cvt.f32.f16 %f349, %rs183;} // end inline asm fma.rn.ftz.f32 %f464, %f457, %f349, %f776; cvt.s16.s8 %rs296, %rs53; shr.s16 %rs297, %rs296, 7; and.b16 %rs298, %rs297, -16; or.b16 %rs299, %rs298, %rs5; cvt.rn.f32.s16 %f465, %rs299; sub.ftz.f32 %f466, %f465, %f9; mul.ftz.f32 %f467, %f77, %f466; // begin inline asm { cvt.f32.f16 %f350, %rs6;} // end inline asm fma.rn.ftz.f32 %f468, %f467, %f350, %f458; // begin inline asm { cvt.f32.f16 %f351, %rs7;} // end inline asm fma.rn.ftz.f32 %f469, %f467, %f351, %f459; // begin inline asm { cvt.f32.f16 %f352, %rs8;} // end inline asm fma.rn.ftz.f32 %f470, %f467, %f352, %f460; // begin inline asm { cvt.f32.f16 %f353, %rs187;} // end inline asm fma.rn.ftz.f32 %f471, %f467, %f353, %f461; // begin inline asm { cvt.f32.f16 %f354, %rs188;} // end inline asm fma.rn.ftz.f32 %f472, %f467, %f354, %f462; // begin inline asm { cvt.f32.f16 %f355, %rs189;} // end inline asm fma.rn.ftz.f32 %f473, %f467, %f355, %f463; // begin inline asm { cvt.f32.f16 %f356, %rs190;} // end inline asm fma.rn.ftz.f32 %f474, %f467, %f356, %f464; cvt.s16.s8 %rs301, %rs54; shr.s16 %rs302, %rs301, 7; and.b16 %rs303, %rs302, -16; or.b16 %rs304, %rs303, %rs9; cvt.rn.f32.s16 %f475, %rs304; sub.ftz.f32 %f476, %f475, %f9; mul.ftz.f32 %f477, %f77, %f476; mov.b32 {%rs191, %rs198}, %r74; // begin inline asm { cvt.f32.f16 %f357, %rs191;} // end inline asm fma.rn.ftz.f32 %f478, %f477, %f357, %f468; mov.b32 {%rs192, %rs199}, %r78; // begin inline asm { cvt.f32.f16 %f358, %rs192;} // end inline asm fma.rn.ftz.f32 %f479, %f477, %f358, %f469; mov.b32 {%rs193, %rs200}, %r84; // begin inline asm { cvt.f32.f16 %f359, %rs193;} // end inline asm fma.rn.ftz.f32 %f480, %f477, %f359, %f470; mov.b32 {%rs194, %rs201}, %r88; // begin inline asm { cvt.f32.f16 %f360, %rs194;} // end inline asm fma.rn.ftz.f32 %f481, %f477, %f360, %f471; mov.b32 {%rs195, %rs202}, %r92; // begin inline asm { cvt.f32.f16 %f361, %rs195;} // end inline asm fma.rn.ftz.f32 %f482, %f477, %f361, %f472; mov.b32 {%rs196, %rs203}, %r96; // begin inline asm { cvt.f32.f16 %f362, %rs196;} // end inline asm fma.rn.ftz.f32 %f483, %f477, %f362, %f473; mov.b32 {%rs197, %rs204}, %r100; // begin inline asm { cvt.f32.f16 %f363, %rs197;} // end inline asm fma.rn.ftz.f32 %f484, %f477, %f363, %f474; shr.s16 %rs305, %rs53, 15; and.b16 %rs306, %rs305, -16; or.b16 %rs307, %rs306, %rs10; cvt.rn.f32.s16 %f485, %rs307; sub.ftz.f32 %f486, %f485, %f9; mul.ftz.f32 %f487, %f77, %f486; // begin inline asm { cvt.f32.f16 %f364, %rs198;} // end inline asm fma.rn.ftz.f32 %f488, %f487, %f364, %f478; // begin inline asm { cvt.f32.f16 %f365, %rs199;} // end inline asm fma.rn.ftz.f32 %f489, %f487, %f365, %f479; // begin inline asm { cvt.f32.f16 %f366, %rs200;} // end inline asm fma.rn.ftz.f32 %f490, %f487, %f366, %f480; // begin inline asm { cvt.f32.f16 %f367, %rs201;} // end inline asm fma.rn.ftz.f32 %f491, %f487, %f367, %f481; // begin inline asm { cvt.f32.f16 %f368, %rs202;} // end inline asm fma.rn.ftz.f32 %f492, %f487, %f368, %f482; // begin inline asm { cvt.f32.f16 %f369, %rs203;} // end inline asm fma.rn.ftz.f32 %f493, %f487, %f369, %f483; // begin inline asm { cvt.f32.f16 %f370, %rs204;} // end inline asm fma.rn.ftz.f32 %f494, %f487, %f370, %f484; shl.b16 %rs309, %rs59, 4; cvt.s16.s8 %rs310, %rs309; shr.s16 %rs311, %rs310, 7; and.b16 %rs312, %rs311, -16; or.b16 %rs313, %rs312, %rs11; cvt.rn.f32.s16 %f495, %rs313; sub.ftz.f32 %f496, %f495, %f9; mul.ftz.f32 %f497, %f77, %f496; mov.b32 {%rs205, %rs212}, %r75; // begin inline asm { cvt.f32.f16 %f371, %rs205;} // end inline asm fma.rn.ftz.f32 %f498, %f497, %f371, %f488; mov.b32 {%rs206, %rs213}, %r79; // begin inline asm { cvt.f32.f16 %f372, %rs206;} // end inline asm fma.rn.ftz.f32 %f499, %f497, %f372, %f489; mov.b32 {%rs207, %rs214}, %r85; // begin inline asm { cvt.f32.f16 %f373, %rs207;} // end inline asm fma.rn.ftz.f32 %f500, %f497, %f373, %f490; mov.b32 {%rs208, %rs215}, %r89; // begin inline asm { cvt.f32.f16 %f374, %rs208;} // end inline asm fma.rn.ftz.f32 %f501, %f497, %f374, %f491; mov.b32 {%rs209, %rs216}, %r93; // begin inline asm { cvt.f32.f16 %f375, %rs209;} // end inline asm fma.rn.ftz.f32 %f502, %f497, %f375, %f492; mov.b32 {%rs210, %rs217}, %r97; // begin inline asm { cvt.f32.f16 %f376, %rs210;} // end inline asm fma.rn.ftz.f32 %f503, %f497, %f376, %f493; mov.b32 {%rs211, %rs218}, %r101; // begin inline asm { cvt.f32.f16 %f377, %rs211;} // end inline asm fma.rn.ftz.f32 %f504, %f497, %f377, %f494; shl.b16 %rs315, %rs60, 4; cvt.s16.s8 %rs316, %rs315; shr.s16 %rs317, %rs316, 7; and.b16 %rs318, %rs317, -16; or.b16 %rs319, %rs318, %rs12; cvt.rn.f32.s16 %f505, %rs319; sub.ftz.f32 %f506, %f505, %f9; mul.ftz.f32 %f507, %f77, %f506; // begin inline asm { cvt.f32.f16 %f378, %rs212;} // end inline asm fma.rn.ftz.f32 %f508, %f507, %f378, %f498; // begin inline asm { cvt.f32.f16 %f379, %rs213;} // end inline asm fma.rn.ftz.f32 %f509, %f507, %f379, %f499; // begin inline asm { cvt.f32.f16 %f380, %rs214;} // end inline asm fma.rn.ftz.f32 %f510, %f507, %f380, %f500; // begin inline asm { cvt.f32.f16 %f381, %rs215;} // end inline asm fma.rn.ftz.f32 %f511, %f507, %f381, %f501; // begin inline asm { cvt.f32.f16 %f382, %rs216;} // end inline asm fma.rn.ftz.f32 %f512, %f507, %f382, %f502; // begin inline asm { cvt.f32.f16 %f383, %rs217;} // end inline asm fma.rn.ftz.f32 %f513, %f507, %f383, %f503; // begin inline asm { cvt.f32.f16 %f384, %rs218;} // end inline asm fma.rn.ftz.f32 %f514, %f507, %f384, %f504; shl.b16 %rs321, %rs61, 4; cvt.s16.s8 %rs322, %rs321; shr.s16 %rs323, %rs322, 7; and.b16 %rs324, %rs323, -16; or.b16 %rs325, %rs324, %rs13; cvt.rn.f32.s16 %f515, %rs325; sub.ftz.f32 %f516, %f515, %f9; mul.ftz.f32 %f517, %f77, %f516; mov.b32 {%rs219, %rs226}, %r76; // begin inline asm { cvt.f32.f16 %f385, %rs219;} // end inline asm fma.rn.ftz.f32 %f518, %f517, %f385, %f508; mov.b32 {%rs220, %rs227}, %r80; // begin inline asm { cvt.f32.f16 %f386, %rs220;} // end inline asm fma.rn.ftz.f32 %f519, %f517, %f386, %f509; mov.b32 {%rs221, %rs228}, %r86; // begin inline asm { cvt.f32.f16 %f387, %rs221;} // end inline asm fma.rn.ftz.f32 %f520, %f517, %f387, %f510; mov.b32 {%rs222, %rs229}, %r90; // begin inline asm { cvt.f32.f16 %f388, %rs222;} // end inline asm fma.rn.ftz.f32 %f521, %f517, %f388, %f511; mov.b32 {%rs223, %rs230}, %r94; // begin inline asm { cvt.f32.f16 %f389, %rs223;} // end inline asm fma.rn.ftz.f32 %f522, %f517, %f389, %f512; mov.b32 {%rs224, %rs231}, %r98; // begin inline asm { cvt.f32.f16 %f390, %rs224;} // end inline asm fma.rn.ftz.f32 %f523, %f517, %f390, %f513; mov.b32 {%rs225, %rs232}, %r102; // begin inline asm { cvt.f32.f16 %f391, %rs225;} // end inline asm fma.rn.ftz.f32 %f524, %f517, %f391, %f514; shl.b16 %rs326, %rs14, 4; cvt.s16.s8 %rs327, %rs326; shr.s16 %rs328, %rs327, 7; and.b16 %rs329, %rs328, -16; or.b16 %rs330, %rs329, %rs14; cvt.rn.f32.s16 %f525, %rs330; sub.ftz.f32 %f526, %f525, %f9; mul.ftz.f32 %f527, %f77, %f526; // begin inline asm { cvt.f32.f16 %f392, %rs226;} // end inline asm fma.rn.ftz.f32 %f528, %f527, %f392, %f518; // begin inline asm { cvt.f32.f16 %f393, %rs227;} // end inline asm fma.rn.ftz.f32 %f529, %f527, %f393, %f519; // begin inline asm { cvt.f32.f16 %f394, %rs228;} // end inline asm fma.rn.ftz.f32 %f530, %f527, %f394, %f520; // begin inline asm { cvt.f32.f16 %f395, %rs229;} // end inline asm fma.rn.ftz.f32 %f531, %f527, %f395, %f521; // begin inline asm { cvt.f32.f16 %f396, %rs230;} // end inline asm fma.rn.ftz.f32 %f532, %f527, %f396, %f522; // begin inline asm { cvt.f32.f16 %f397, %rs231;} // end inline asm fma.rn.ftz.f32 %f533, %f527, %f397, %f523; // begin inline asm { cvt.f32.f16 %f398, %rs232;} // end inline asm fma.rn.ftz.f32 %f534, %f527, %f398, %f524; ld.global.v4.u32 {%r173, %r174, %r175, %r176}, [%rd5+16]; ld.global.v4.u32 {%r181, %r182, %r183, %r184}, [%rd6+16]; ld.global.v4.u32 {%r189, %r190, %r191, %r192}, [%rd7]; ld.global.v4.u32 {%r197, %r198, %r199, %r200}, [%rd8]; ld.global.v4.u32 {%r205, %r206, %r207, %r208}, [%rd9]; ld.global.v4.u32 {%r213, %r214, %r215, %r216}, [%rd10]; ld.global.v4.u32 {%r221, %r222, %r223, %r224}, [%rd11]; shl.b16 %rs332, %rs62, 4; cvt.s16.s8 %rs333, %rs332; shr.s16 %rs334, %rs333, 7; and.b16 %rs335, %rs334, -16; or.b16 %rs336, %rs335, %rs15; cvt.rn.f32.s16 %f535, %rs336; sub.ftz.f32 %f536, %f535, %f9; mul.ftz.f32 %f537, %f77, %f536; mov.b32 {%rs233, %rs240}, %r173; // begin inline asm { cvt.f32.f16 %f399, %rs233;} // end inline asm fma.rn.ftz.f32 %f538, %f537, %f399, %f528; mov.b32 {%rs234, %rs241}, %r181; // begin inline asm { cvt.f32.f16 %f400, %rs234;} // end inline asm fma.rn.ftz.f32 %f539, %f537, %f400, %f529; mov.b32 {%rs235, %rs242}, %r189; // begin inline asm { cvt.f32.f16 %f401, %rs235;} // end inline asm fma.rn.ftz.f32 %f540, %f537, %f401, %f530; mov.b32 {%rs236, %rs243}, %r197; // begin inline asm { cvt.f32.f16 %f402, %rs236;} // end inline asm fma.rn.ftz.f32 %f541, %f537, %f402, %f531; mov.b32 {%rs237, %rs244}, %r205; // begin inline asm { cvt.f32.f16 %f403, %rs237;} // end inline asm fma.rn.ftz.f32 %f542, %f537, %f403, %f532; mov.b32 {%rs238, %rs245}, %r213; // begin inline asm { cvt.f32.f16 %f404, %rs238;} // end inline asm fma.rn.ftz.f32 %f543, %f537, %f404, %f533; mov.b32 {%rs239, %rs246}, %r221; // begin inline asm { cvt.f32.f16 %f405, %rs239;} // end inline asm fma.rn.ftz.f32 %f544, %f537, %f405, %f534; shl.b16 %rs337, %rs16, 4; cvt.s16.s8 %rs338, %rs337; shr.s16 %rs339, %rs338, 7; and.b16 %rs340, %rs339, -16; or.b16 %rs341, %rs340, %rs17; cvt.rn.f32.s16 %f545, %rs341; sub.ftz.f32 %f546, %f545, %f9; mul.ftz.f32 %f547, %f77, %f546; // begin inline asm { cvt.f32.f16 %f406, %rs240;} // end inline asm fma.rn.ftz.f32 %f548, %f547, %f406, %f538; // begin inline asm { cvt.f32.f16 %f407, %rs241;} // end inline asm fma.rn.ftz.f32 %f549, %f547, %f407, %f539; // begin inline asm { cvt.f32.f16 %f408, %rs242;} // end inline asm fma.rn.ftz.f32 %f550, %f547, %f408, %f540; // begin inline asm { cvt.f32.f16 %f409, %rs243;} // end inline asm fma.rn.ftz.f32 %f551, %f547, %f409, %f541; // begin inline asm { cvt.f32.f16 %f410, %rs244;} // end inline asm fma.rn.ftz.f32 %f552, %f547, %f410, %f542; // begin inline asm { cvt.f32.f16 %f411, %rs245;} // end inline asm fma.rn.ftz.f32 %f553, %f547, %f411, %f543; // begin inline asm { cvt.f32.f16 %f412, %rs246;} // end inline asm fma.rn.ftz.f32 %f554, %f547, %f412, %f544; shl.b16 %rs342, %rs18, 4; cvt.s16.s8 %rs343, %rs342; shr.s16 %rs344, %rs343, 7; and.b16 %rs345, %rs344, -16; or.b16 %rs346, %rs345, %rs19; cvt.rn.f32.s16 %f555, %rs346; sub.ftz.f32 %f556, %f555, %f9; mul.ftz.f32 %f557, %f77, %f556; mov.b32 {%rs247, %rs254}, %r174; // begin inline asm { cvt.f32.f16 %f413, %rs247;} // end inline asm fma.rn.ftz.f32 %f558, %f557, %f413, %f548; mov.b32 {%rs248, %rs255}, %r182; // begin inline asm { cvt.f32.f16 %f414, %rs248;} // end inline asm fma.rn.ftz.f32 %f559, %f557, %f414, %f549; mov.b32 {%rs249, %rs256}, %r190; // begin inline asm { cvt.f32.f16 %f415, %rs249;} // end inline asm fma.rn.ftz.f32 %f560, %f557, %f415, %f550; mov.b32 {%rs250, %rs257}, %r198; // begin inline asm { cvt.f32.f16 %f416, %rs250;} // end inline asm fma.rn.ftz.f32 %f561, %f557, %f416, %f551; mov.b32 {%rs251, %rs258}, %r206; // begin inline asm { cvt.f32.f16 %f417, %rs251;} // end inline asm fma.rn.ftz.f32 %f562, %f557, %f417, %f552; mov.b32 {%rs252, %rs259}, %r214; // begin inline asm { cvt.f32.f16 %f418, %rs252;} // end inline asm fma.rn.ftz.f32 %f563, %f557, %f418, %f553; mov.b32 {%rs253, %rs260}, %r222; // begin inline asm { cvt.f32.f16 %f419, %rs253;} // end inline asm fma.rn.ftz.f32 %f564, %f557, %f419, %f554; shl.b16 %rs347, %rs20, 4; cvt.s16.s8 %rs348, %rs347; shr.s16 %rs349, %rs348, 7; and.b16 %rs350, %rs349, -16; or.b16 %rs351, %rs350, %rs21; cvt.rn.f32.s16 %f565, %rs351; sub.ftz.f32 %f566, %f565, %f9; mul.ftz.f32 %f567, %f77, %f566; // begin inline asm { cvt.f32.f16 %f420, %rs254;} // end inline asm fma.rn.ftz.f32 %f568, %f567, %f420, %f558; // begin inline asm { cvt.f32.f16 %f421, %rs255;} // end inline asm fma.rn.ftz.f32 %f569, %f567, %f421, %f559; // begin inline asm { cvt.f32.f16 %f422, %rs256;} // end inline asm fma.rn.ftz.f32 %f570, %f567, %f422, %f560; // begin inline asm { cvt.f32.f16 %f423, %rs257;} // end inline asm fma.rn.ftz.f32 %f571, %f567, %f423, %f561; // begin inline asm { cvt.f32.f16 %f424, %rs258;} // end inline asm fma.rn.ftz.f32 %f572, %f567, %f424, %f562; // begin inline asm { cvt.f32.f16 %f425, %rs259;} // end inline asm fma.rn.ftz.f32 %f573, %f567, %f425, %f563; // begin inline asm { cvt.f32.f16 %f426, %rs260;} // end inline asm fma.rn.ftz.f32 %f574, %f567, %f426, %f564; shl.b16 %rs352, %rs22, 4; cvt.s16.s8 %rs353, %rs352; shr.s16 %rs354, %rs353, 7; and.b16 %rs355, %rs354, -16; or.b16 %rs356, %rs355, %rs23; cvt.rn.f32.s16 %f575, %rs356; sub.ftz.f32 %f576, %f575, %f9; mul.ftz.f32 %f577, %f77, %f576; mov.b32 {%rs261, %rs268}, %r175; // begin inline asm { cvt.f32.f16 %f427, %rs261;} // end inline asm fma.rn.ftz.f32 %f578, %f577, %f427, %f568; mov.b32 {%rs262, %rs269}, %r183; // begin inline asm { cvt.f32.f16 %f428, %rs262;} // end inline asm fma.rn.ftz.f32 %f579, %f577, %f428, %f569; mov.b32 {%rs263, %rs270}, %r191; // begin inline asm { cvt.f32.f16 %f429, %rs263;} // end inline asm fma.rn.ftz.f32 %f580, %f577, %f429, %f570; mov.b32 {%rs264, %rs271}, %r199; // begin inline asm { cvt.f32.f16 %f430, %rs264;} // end inline asm fma.rn.ftz.f32 %f581, %f577, %f430, %f571; mov.b32 {%rs265, %rs272}, %r207; // begin inline asm { cvt.f32.f16 %f431, %rs265;} // end inline asm fma.rn.ftz.f32 %f582, %f577, %f431, %f572; mov.b32 {%rs266, %rs273}, %r215; // begin inline asm { cvt.f32.f16 %f432, %rs266;} // end inline asm fma.rn.ftz.f32 %f583, %f577, %f432, %f573; mov.b32 {%rs267, %rs274}, %r223; // begin inline asm { cvt.f32.f16 %f433, %rs267;} // end inline asm fma.rn.ftz.f32 %f584, %f577, %f433, %f574; shl.b16 %rs357, %rs24, 4; cvt.s16.s8 %rs358, %rs357; shr.s16 %rs359, %rs358, 7; and.b16 %rs360, %rs359, -16; or.b16 %rs361, %rs360, %rs25; cvt.rn.f32.s16 %f585, %rs361; sub.ftz.f32 %f586, %f585, %f9; mul.ftz.f32 %f587, %f77, %f586; // begin inline asm { cvt.f32.f16 %f434, %rs268;} // end inline asm fma.rn.ftz.f32 %f588, %f587, %f434, %f578; // begin inline asm { cvt.f32.f16 %f435, %rs269;} // end inline asm fma.rn.ftz.f32 %f589, %f587, %f435, %f579; // begin inline asm { cvt.f32.f16 %f436, %rs270;} // end inline asm fma.rn.ftz.f32 %f590, %f587, %f436, %f580; // begin inline asm { cvt.f32.f16 %f437, %rs271;} // end inline asm fma.rn.ftz.f32 %f591, %f587, %f437, %f581; // begin inline asm { cvt.f32.f16 %f438, %rs272;} // end inline asm fma.rn.ftz.f32 %f592, %f587, %f438, %f582; // begin inline asm { cvt.f32.f16 %f439, %rs273;} // end inline asm fma.rn.ftz.f32 %f593, %f587, %f439, %f583; // begin inline asm { cvt.f32.f16 %f440, %rs274;} // end inline asm fma.rn.ftz.f32 %f594, %f587, %f440, %f584; shl.b16 %rs362, %rs26, 4; cvt.s16.s8 %rs363, %rs362; shr.s16 %rs364, %rs363, 7; and.b16 %rs365, %rs364, -16; or.b16 %rs366, %rs365, %rs27; cvt.rn.f32.s16 %f595, %rs366; sub.ftz.f32 %f596, %f595, %f9; mul.ftz.f32 %f597, %f77, %f596; mov.b32 {%rs275, %rs282}, %r176; // begin inline asm { cvt.f32.f16 %f441, %rs275;} // end inline asm fma.rn.ftz.f32 %f598, %f597, %f441, %f588; mov.b32 {%rs276, %rs283}, %r184; // begin inline asm { cvt.f32.f16 %f442, %rs276;} // end inline asm fma.rn.ftz.f32 %f599, %f597, %f442, %f589; mov.b32 {%rs277, %rs284}, %r192; // begin inline asm { cvt.f32.f16 %f443, %rs277;} // end inline asm fma.rn.ftz.f32 %f600, %f597, %f443, %f590; mov.b32 {%rs278, %rs285}, %r200; // begin inline asm { cvt.f32.f16 %f444, %rs278;} // end inline asm fma.rn.ftz.f32 %f601, %f597, %f444, %f591; mov.b32 {%rs279, %rs286}, %r208; // begin inline asm { cvt.f32.f16 %f445, %rs279;} // end inline asm fma.rn.ftz.f32 %f602, %f597, %f445, %f592; mov.b32 {%rs280, %rs287}, %r216; // begin inline asm { cvt.f32.f16 %f446, %rs280;} // end inline asm fma.rn.ftz.f32 %f603, %f597, %f446, %f593; mov.b32 {%rs281, %rs288}, %r224; // begin inline asm { cvt.f32.f16 %f447, %rs281;} // end inline asm fma.rn.ftz.f32 %f604, %f597, %f447, %f594; shl.b16 %rs367, %rs28, 4; cvt.s16.s8 %rs368, %rs367; shr.s16 %rs369, %rs368, 7; and.b16 %rs370, %rs369, -16; or.b16 %rs371, %rs370, %rs28; cvt.rn.f32.s16 %f605, %rs371; sub.ftz.f32 %f606, %f605, %f9; mul.ftz.f32 %f607, %f77, %f606; // begin inline asm { cvt.f32.f16 %f448, %rs282;} // end inline asm fma.rn.ftz.f32 %f782, %f607, %f448, %f598; // begin inline asm { cvt.f32.f16 %f449, %rs283;} // end inline asm fma.rn.ftz.f32 %f781, %f607, %f449, %f599; // begin inline asm { cvt.f32.f16 %f450, %rs284;} // end inline asm fma.rn.ftz.f32 %f780, %f607, %f450, %f600; // begin inline asm { cvt.f32.f16 %f451, %rs285;} // end inline asm fma.rn.ftz.f32 %f779, %f607, %f451, %f601; // begin inline asm { cvt.f32.f16 %f452, %rs286;} // end inline asm fma.rn.ftz.f32 %f778, %f607, %f452, %f602; // begin inline asm { cvt.f32.f16 %f453, %rs287;} // end inline asm fma.rn.ftz.f32 %f777, %f607, %f453, %f603; // begin inline asm { cvt.f32.f16 %f454, %rs288;} // end inline asm fma.rn.ftz.f32 %f776, %f607, %f454, %f604; $L__BB0_8: add.s32 %r370, %r370, 4; shl.b32 %r229, %r370, 5; add.s32 %r369, %r229, %r57; shl.b32 %r368, %r369, 1; setp.lt.u32 %p7, %r368, %r54; @%p7 bra $L__BB0_2; $L__BB0_9: mov.u32 %r363, %tid.y; mov.u32 %r362, %tid.x; shl.b32 %r361, %r363, 5; add.s32 %r360, %r361, %r362; shl.b32 %r231, %r360, 2; mov.u32 %r232, _ZZ9gemv_int4ILi4ELi64ELi7EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage; add.s32 %r233, %r232, %r231; setp.lt.u32 %p8, %r360, 32; @%p8 bra $L__BB0_11; add.s32 %r353, %r233, -112; st.shared.f32 [%r353], %f782; $L__BB0_11: mov.u32 %r367, %tid.y; mov.u32 %r366, %tid.x; shl.b32 %r365, %r367, 5; add.s32 %r364, %r365, %r366; setp.gt.u32 %p9, %r364, 31; bar.sync 0; mad.lo.s32 %r47, %r364, 12, %r232; @%p9 bra $L__BB0_13; mov.u32 %r248, 16; ld.shared.f32 %f623, [%r47+16]; add.ftz.f32 %f624, %f782, %f623; ld.shared.f32 %f625, [%r47+20]; add.ftz.f32 %f626, %f624, %f625; ld.shared.f32 %f627, [%r47+24]; add.ftz.f32 %f610, %f626, %f627; mov.u32 %r236, 1; mov.u32 %r249, 31; mov.u32 %r250, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f610, %r236, %r249, %r250; @p add.f32 r0, r0, %f610; mov.f32 %f608, r0;} // end inline asm mov.u32 %r239, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f608, %r239, %r249, %r250; @p add.f32 r0, r0, %f608; mov.f32 %f611, r0;} // end inline asm mov.u32 %r242, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f611, %r242, %r249, %r250; @p add.f32 r0, r0, %f611; mov.f32 %f614, r0;} // end inline asm mov.u32 %r245, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f614, %r245, %r249, %r250; @p add.f32 r0, r0, %f614; mov.f32 %f617, r0;} // end inline asm // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f617, %r248, %r249, %r250; @p add.f32 r0, r0, %f617; mov.f32 %f782, r0;} // end inline asm $L__BB0_13: @%p8 bra $L__BB0_15; add.s32 %r354, %r233, -112; st.shared.f32 [%r354+640], %f781; $L__BB0_15: bar.sync 0; @%p9 bra $L__BB0_17; ld.shared.f32 %f643, [%r47+656]; add.ftz.f32 %f644, %f781, %f643; ld.shared.f32 %f645, [%r47+660]; add.ftz.f32 %f646, %f644, %f645; ld.shared.f32 %f647, [%r47+664]; add.ftz.f32 %f630, %f646, %f647; mov.u32 %r252, 1; mov.u32 %r265, 31; mov.u32 %r266, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f630, %r252, %r265, %r266; @p add.f32 r0, r0, %f630; mov.f32 %f628, r0;} // end inline asm mov.u32 %r255, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f628, %r255, %r265, %r266; @p add.f32 r0, r0, %f628; mov.f32 %f631, r0;} // end inline asm mov.u32 %r258, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f631, %r258, %r265, %r266; @p add.f32 r0, r0, %f631; mov.f32 %f634, r0;} // end inline asm mov.u32 %r261, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f634, %r261, %r265, %r266; @p add.f32 r0, r0, %f634; mov.f32 %f637, r0;} // end inline asm mov.u32 %r264, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f637, %r264, %r265, %r266; @p add.f32 r0, r0, %f637; mov.f32 %f781, r0;} // end inline asm $L__BB0_17: @%p8 bra $L__BB0_19; add.s32 %r355, %r233, -112; st.shared.f32 [%r355+1280], %f780; $L__BB0_19: bar.sync 0; @%p9 bra $L__BB0_21; ld.shared.f32 %f663, [%r47+1296]; add.ftz.f32 %f664, %f780, %f663; ld.shared.f32 %f665, [%r47+1300]; add.ftz.f32 %f666, %f664, %f665; ld.shared.f32 %f667, [%r47+1304]; add.ftz.f32 %f650, %f666, %f667; mov.u32 %r268, 1; mov.u32 %r281, 31; mov.u32 %r282, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f650, %r268, %r281, %r282; @p add.f32 r0, r0, %f650; mov.f32 %f648, r0;} // end inline asm mov.u32 %r271, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f648, %r271, %r281, %r282; @p add.f32 r0, r0, %f648; mov.f32 %f651, r0;} // end inline asm mov.u32 %r274, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f651, %r274, %r281, %r282; @p add.f32 r0, r0, %f651; mov.f32 %f654, r0;} // end inline asm mov.u32 %r277, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f654, %r277, %r281, %r282; @p add.f32 r0, r0, %f654; mov.f32 %f657, r0;} // end inline asm mov.u32 %r280, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f657, %r280, %r281, %r282; @p add.f32 r0, r0, %f657; mov.f32 %f780, r0;} // end inline asm $L__BB0_21: @%p8 bra $L__BB0_23; add.s32 %r356, %r233, -112; st.shared.f32 [%r356+1920], %f779; $L__BB0_23: bar.sync 0; @%p9 bra $L__BB0_25; ld.shared.f32 %f683, [%r47+1936]; add.ftz.f32 %f684, %f779, %f683; ld.shared.f32 %f685, [%r47+1940]; add.ftz.f32 %f686, %f684, %f685; ld.shared.f32 %f687, [%r47+1944]; add.ftz.f32 %f670, %f686, %f687; mov.u32 %r284, 1; mov.u32 %r297, 31; mov.u32 %r298, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f670, %r284, %r297, %r298; @p add.f32 r0, r0, %f670; mov.f32 %f668, r0;} // end inline asm mov.u32 %r287, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f668, %r287, %r297, %r298; @p add.f32 r0, r0, %f668; mov.f32 %f671, r0;} // end inline asm mov.u32 %r290, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f671, %r290, %r297, %r298; @p add.f32 r0, r0, %f671; mov.f32 %f674, r0;} // end inline asm mov.u32 %r293, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f674, %r293, %r297, %r298; @p add.f32 r0, r0, %f674; mov.f32 %f677, r0;} // end inline asm mov.u32 %r296, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f677, %r296, %r297, %r298; @p add.f32 r0, r0, %f677; mov.f32 %f779, r0;} // end inline asm $L__BB0_25: @%p8 bra $L__BB0_27; add.s32 %r357, %r233, -112; st.shared.f32 [%r357+2560], %f778; $L__BB0_27: bar.sync 0; @%p9 bra $L__BB0_29; ld.shared.f32 %f703, [%r47+2576]; add.ftz.f32 %f704, %f778, %f703; ld.shared.f32 %f705, [%r47+2580]; add.ftz.f32 %f706, %f704, %f705; ld.shared.f32 %f707, [%r47+2584]; add.ftz.f32 %f690, %f706, %f707; mov.u32 %r300, 1; mov.u32 %r313, 31; mov.u32 %r314, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f690, %r300, %r313, %r314; @p add.f32 r0, r0, %f690; mov.f32 %f688, r0;} // end inline asm mov.u32 %r303, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f688, %r303, %r313, %r314; @p add.f32 r0, r0, %f688; mov.f32 %f691, r0;} // end inline asm mov.u32 %r306, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f691, %r306, %r313, %r314; @p add.f32 r0, r0, %f691; mov.f32 %f694, r0;} // end inline asm mov.u32 %r309, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f694, %r309, %r313, %r314; @p add.f32 r0, r0, %f694; mov.f32 %f697, r0;} // end inline asm mov.u32 %r312, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f697, %r312, %r313, %r314; @p add.f32 r0, r0, %f697; mov.f32 %f778, r0;} // end inline asm $L__BB0_29: @%p8 bra $L__BB0_31; add.s32 %r358, %r233, -112; st.shared.f32 [%r358+3200], %f777; $L__BB0_31: bar.sync 0; @%p9 bra $L__BB0_33; ld.shared.f32 %f723, [%r47+3216]; add.ftz.f32 %f724, %f777, %f723; ld.shared.f32 %f725, [%r47+3220]; add.ftz.f32 %f726, %f724, %f725; ld.shared.f32 %f727, [%r47+3224]; add.ftz.f32 %f710, %f726, %f727; mov.u32 %r316, 1; mov.u32 %r329, 31; mov.u32 %r330, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f710, %r316, %r329, %r330; @p add.f32 r0, r0, %f710; mov.f32 %f708, r0;} // end inline asm mov.u32 %r319, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f708, %r319, %r329, %r330; @p add.f32 r0, r0, %f708; mov.f32 %f711, r0;} // end inline asm mov.u32 %r322, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f711, %r322, %r329, %r330; @p add.f32 r0, r0, %f711; mov.f32 %f714, r0;} // end inline asm mov.u32 %r325, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f714, %r325, %r329, %r330; @p add.f32 r0, r0, %f714; mov.f32 %f717, r0;} // end inline asm mov.u32 %r328, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f717, %r328, %r329, %r330; @p add.f32 r0, r0, %f717; mov.f32 %f777, r0;} // end inline asm $L__BB0_33: @%p8 bra $L__BB0_35; add.s32 %r359, %r233, -112; st.shared.f32 [%r359+3840], %f776; $L__BB0_35: bar.sync 0; @%p9 bra $L__BB0_37; ld.shared.f32 %f743, [%r47+3856]; add.ftz.f32 %f744, %f776, %f743; ld.shared.f32 %f745, [%r47+3860]; add.ftz.f32 %f746, %f744, %f745; ld.shared.f32 %f747, [%r47+3864]; add.ftz.f32 %f730, %f746, %f747; mov.u32 %r332, 1; mov.u32 %r345, 31; mov.u32 %r346, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f730, %r332, %r345, %r346; @p add.f32 r0, r0, %f730; mov.f32 %f728, r0;} // end inline asm mov.u32 %r335, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f728, %r335, %r345, %r346; @p add.f32 r0, r0, %f728; mov.f32 %f731, r0;} // end inline asm mov.u32 %r338, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f731, %r338, %r345, %r346; @p add.f32 r0, r0, %f731; mov.f32 %f734, r0;} // end inline asm mov.u32 %r341, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f734, %r341, %r345, %r346; @p add.f32 r0, r0, %f734; mov.f32 %f737, r0;} // end inline asm mov.u32 %r344, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f737, %r344, %r345, %r346; @p add.f32 r0, r0, %f737; mov.f32 %f776, r0;} // end inline asm $L__BB0_37: mov.u32 %r347, %tid.y; or.b32 %r349, %r57, %r347; setp.ne.s32 %p22, %r349, 0; @%p22 bra $L__BB0_41; ld.param.u64 %rd67, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+8]; ld.param.u64 %rd66, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0]; cvta.to.global.u64 %rd42, %rd66; setp.eq.s64 %p23, %rd67, 0; mul.ftz.f32 %f52, %f68, %f782; mov.u32 %r350, %ctaid.x; cvt.s64.s32 %rd13, %r350; mul.wide.s32 %rd43, %r350, 2; add.s64 %rd14, %rd42, %rd43; mul.ftz.f32 %f53, %f68, %f781; add.s32 %r351, %r53, %r350; cvt.s64.s32 %rd15, %r53; mul.wide.s32 %rd44, %r53, 2; add.s64 %rd16, %rd14, %rd44; mul.ftz.f32 %f54, %f68, %f780; add.s32 %r352, %r351, %r53; cvt.s64.s32 %rd17, %r352; mul.wide.s32 %rd45, %r352, 2; add.s64 %rd19, %rd42, %rd45; mul.ftz.f32 %f55, %f68, %f779; mul.ftz.f32 %f56, %f68, %f778; mul.ftz.f32 %f57, %f68, %f777; mul.ftz.f32 %f58, %f68, %f776; @%p23 bra $L__BB0_40; ld.param.u64 %rd68, [_Z27dequant_gemv_group64_batch723DequantGemvKernelParams_param_0+8]; cvta.to.global.u64 %rd46, %rd68; shl.b64 %rd47, %rd13, 1; add.s64 %rd48, %rd46, %rd47; ld.global.u16 %rs372, [%rd48]; // begin inline asm { cvt.f32.f16 %f748, %rs372;} // end inline asm fma.rn.ftz.f32 %f749, %f69, %f748, %f52; // begin inline asm { cvt.rn.f16.f32 %rs373, %f749;} // end inline asm st.global.u16 [%rd14], %rs373; shl.b64 %rd49, %rd15, 1; add.s64 %rd50, %rd48, %rd49; ld.global.u16 %rs374, [%rd50]; // begin inline asm { cvt.f32.f16 %f750, %rs374;} // end inline asm fma.rn.ftz.f32 %f751, %f69, %f750, %f53; // begin inline asm { cvt.rn.f16.f32 %rs375, %f751;} // end inline asm st.global.u16 [%rd16], %rs375; shl.b64 %rd51, %rd17, 1; add.s64 %rd52, %rd46, %rd51; ld.global.u16 %rs376, [%rd52]; // begin inline asm { cvt.f32.f16 %f752, %rs376;} // end inline asm fma.rn.ftz.f32 %f753, %f69, %f752, %f54; // begin inline asm { cvt.rn.f16.f32 %rs377, %f753;} // end inline asm st.global.u16 [%rd19], %rs377; add.s64 %rd53, %rd52, %rd49; ld.global.u16 %rs378, [%rd53]; // begin inline asm { cvt.f32.f16 %f754, %rs378;} // end inline asm fma.rn.ftz.f32 %f755, %f69, %f754, %f55; // begin inline asm { cvt.rn.f16.f32 %rs379, %f755;} // end inline asm add.s64 %rd54, %rd19, %rd49; st.global.u16 [%rd54], %rs379; add.s64 %rd55, %rd53, %rd49; ld.global.u16 %rs380, [%rd55]; // begin inline asm { cvt.f32.f16 %f756, %rs380;} // end inline asm fma.rn.ftz.f32 %f757, %f69, %f756, %f56; // begin inline asm { cvt.rn.f16.f32 %rs381, %f757;} // end inline asm add.s64 %rd56, %rd54, %rd49; st.global.u16 [%rd56], %rs381; add.s64 %rd57, %rd55, %rd49; ld.global.u16 %rs382, [%rd57]; // begin inline asm { cvt.f32.f16 %f758, %rs382;} // end inline asm fma.rn.ftz.f32 %f759, %f69, %f758, %f57; // begin inline asm { cvt.rn.f16.f32 %rs383, %f759;} // end inline asm add.s64 %rd58, %rd56, %rd49; st.global.u16 [%rd58], %rs383; add.s64 %rd59, %rd57, %rd49; ld.global.u16 %rs384, [%rd59]; // begin inline asm { cvt.f32.f16 %f760, %rs384;} // end inline asm fma.rn.ftz.f32 %f761, %f69, %f760, %f58; // begin inline asm { cvt.rn.f16.f32 %rs385, %f761;} // end inline asm add.s64 %rd60, %rd58, %rd49; st.global.u16 [%rd60], %rs385; bra.uni $L__BB0_41; $L__BB0_40: // begin inline asm { cvt.rn.f16.f32 %rs386, %f52;} // end inline asm st.global.u16 [%rd14], %rs386; // begin inline asm { cvt.rn.f16.f32 %rs387, %f53;} // end inline asm st.global.u16 [%rd16], %rs387; // begin inline asm { cvt.rn.f16.f32 %rs388, %f54;} // end inline asm st.global.u16 [%rd19], %rs388; // begin inline asm { cvt.rn.f16.f32 %rs389, %f55;} // end inline asm shl.b64 %rd61, %rd15, 1; add.s64 %rd62, %rd19, %rd61; st.global.u16 [%rd62], %rs389; // begin inline asm { cvt.rn.f16.f32 %rs390, %f56;} // end inline asm add.s64 %rd63, %rd62, %rd61; st.global.u16 [%rd63], %rs390; // begin inline asm { cvt.rn.f16.f32 %rs391, %f57;} // end inline asm add.s64 %rd64, %rd63, %rd61; st.global.u16 [%rd64], %rs391; // begin inline asm { cvt.rn.f16.f32 %rs392, %f58;} // end inline asm add.s64 %rd65, %rd64, %rd61; st.global.u16 [%rd65], %rs392; $L__BB0_41: ret; } // .globl _ZN3cub11EmptyKernelIvEEvv .visible .entry _ZN3cub11EmptyKernelIvEEvv() { ret; }