rams_param_0+56]; ld.param.v2.f32 {%f50, %f51}, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+64]; ld.param.v4.u8 {%rs72, %rs73, %rs74, %rs75}, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+72]; ld.param.u64 %rd23, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+40]; ld.param.u64 %rd22, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+32]; ld.param.u64 %rd21, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+24]; ld.param.u64 %rd20, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+16]; mov.u32 %r290, %tid.y; shl.b32 %r48, %r290, 5; mov.u32 %r49, %tid.x; add.s32 %r289, %r48, %r49; shl.b32 %r288, %r289, 1; setp.ge.u32 %p1, %r288, %r46; mov.f32 %f584, 0f00000000; mov.f32 %f585, %f584; mov.f32 %f586, %f584; mov.f32 %f587, %f584; mov.f32 %f588, %f584; @%p1 bra $L__BB0_9; cvta.to.global.u64 %rd2, %rd20; mov.u32 %r50, %ctaid.x; mul.lo.s32 %r6, %r47, %r50; shl.b16 %rs2, %rs72, 3; cvta.to.global.u64 %rd3, %rd21; cvta.to.global.u64 %rd4, %rd23; $L__BB0_2: mad.lo.s32 %r52, %r46, %r50, %r288; mul.wide.u32 %rd24, %r52, 4; add.s64 %rd25, %rd3, %rd24; ld.global.v2.u32 {%r53, %r54}, [%rd25]; shr.u32 %r56, %r49, 2; shl.b32 %r57, %r290, 3; add.s32 %r12, %r57, %r56; add.s32 %r13, %r12, %r6; mul.wide.s32 %rd26, %r13, 2; add.s64 %rd27, %rd4, %rd26; ld.global.u16 %rs80, [%rd27]; // begin inline asm { cvt.f32.f16 %f57, %rs80;} // end inline asm setp.eq.s64 %p2, %rd22, 0; mov.u16 %rs350, %rs2; @%p2 bra $L__BB0_4; shr.u32 %r58, %r13, 31; add.s32 %r59, %r13, %r58; shr.s32 %r60, %r59, 1; cvt.s64.s32 %rd28, %r60; cvta.to.global.u64 %rd29, %rd22; add.s64 %rd30, %rd29, %rd28; ld.global.u8 %r61, [%rd30]; shl.b32 %r62, %r12, 2; and.b32 %r63, %r62, 4; shr.u32 %r64, %r61, %r63; cvt.u16.u32 %rs81, %r64; and.b16 %rs350, %rs81, 15; $L__BB0_4: shl.b32 %r14, %r289, 4; setp.ge.s32 %p3, %r14, %r44; @%p3 bra $L__BB0_8; setp.eq.s16 %p4, %rs72, 0; shr.u16 %rs83, %rs350, 3; and.b16 %rs84, %rs83, 1; setp.eq.b16 %p5, %rs84, 1; and.pred %p6, %p4, %p5; selp.b16 %rs85, -16, 0, %p6; or.b16 %rs86, %rs85, %rs350; cvt.s16.s8 %rs87, %rs86; cvt.rn.f32.s16 %f7, %rs87; mul.wide.s32 %rd31, %r14, 2; add.s64 %rd5, %rd2, %rd31; ld.global.v4.u32 {%r65, %r66, %r67, %r68}, [%rd5]; mul.wide.s32 %rd32, %r44, 2; add.s64 %rd6, %rd5, %rd32; ld.global.v4.u32 {%r69, %r70, %r71, %r72}, [%rd6]; add.s32 %r73, %r14, %r44; add.s32 %r74, %r73, %r44; shl.b32 %r75, %r44, 1; mul.wide.s32 %rd33, %r75, 2; add.s64 %rd34, %rd5, %rd33; ld.global.v4.u32 {%r76, %r77, %r78, %r79}, [%rd34]; add.s32 %r80, %r74, %r44; mul.wide.s32 %rd35, %r80, 2; add.s64 %rd36, %rd2, %rd35; ld.global.v4.u32 {%r81, %r82, %r83, %r84}, [%rd36]; add.s64 %rd37, %rd34, %rd33; ld.global.v4.u32 {%r85, %r86, %r87, %r88}, [%rd37]; cvt.u16.u32 %rs88, %r53; and.b16 %rs5, %rs88, 15; mov.b32 {%rs6, %rs13}, %r65; mov.b32 {%rs8, %rs14}, %r69; mov.b32 {%rs9, %rs15}, %r76; mov.b32 {%rs10, %rs16}, %r81; mov.b32 {%rs11, %rs17}, %r85; shr.u16 %rs89, %rs88, 4; and.b16 %rs12, %rs89, 15; shr.u16 %rs90, %rs88, 8; and.b16 %rs18, %rs90, 15; mov.b32 {%rs19, %rs25}, %r66; mov.b32 {%rs20, %rs26}, %r70; mov.b32 {%rs21, %rs27}, %r77; mov.b32 {%rs22, %rs28}, %r82; mov.b32 {%rs23, %rs29}, %r86; shr.u16 %rs24, %rs88, 12; shr.u32 %r89, %r53, 16; cvt.u16.u32 %rs91, %r89; and.b16 %rs30, %rs91, 15; mov.b32 {%rs31, %rs42}, %r67; mov.b32 {%rs33, %rs43}, %r71; mov.b32 {%rs35, %rs44}, %r78; mov.b32 {%rs37, %rs45}, %r83; mov.b32 {%rs39, %rs46}, %r87; shr.u32 %r90, %r53, 20; cvt.u16.u32 %rs92, %r90; and.b16 %rs41, %rs92, 15; shr.u32 %r91, %r53, 24; cvt.u16.u32 %rs93, %r91; and.b16 %rs47, %rs93, 15; shr.u32 %r92, %r53, 28; cvt.u16.u32 %rs48, %r92; cvt.u16.u32 %rs49, %r54; and.b16 %rs50, %rs49, 15; shr.u32 %r93, %r54, 4; cvt.u16.u32 %rs51, %r93; and.b16 %rs52, %rs51, 15; shr.u32 %r94, %r54, 8; cvt.u16.u32 %rs53, %r94; and.b16 %rs54, %rs53, 15; shr.u32 %r95, %r54, 12; cvt.u16.u32 %rs55, %r95; and.b16 %rs56, %rs55, 15; shr.u32 %r96, %r54, 16; cvt.u16.u32 %rs57, %r96; and.b16 %rs58, %rs57, 15; shr.u32 %r97, %r54, 20; cvt.u16.u32 %rs59, %r97; and.b16 %rs60, %rs59, 15; shr.u32 %r98, %r54, 24; cvt.u16.u32 %rs61, %r98; and.b16 %rs62, %rs61, 15; shr.u32 %r99, %r54, 28; cvt.u16.u32 %rs63, %r99; add.s64 %rd38, %rd6, %rd32; add.s64 %rd7, %rd38, 16; add.s64 %rd8, %rd7, %rd32; add.s64 %rd9, %rd8, %rd32; @%p4 bra $L__BB0_7; cvt.rn.f32.s16 %f138, %rs5; sub.ftz.f32 %f139, %f138, %f7; mul.ftz.f32 %f140, %f57, %f139; // begin inline asm { cvt.f32.f16 %f58, %rs6;} // end inline asm fma.rn.ftz.f32 %f141, %f140, %f58, %f588; // begin inline asm { cvt.f32.f16 %f59, %rs8;} // end inline asm fma.rn.ftz.f32 %f142, %f140, %f59, %f587; // begin inline asm { cvt.f32.f16 %f60, %rs9;} // end inline asm fma.rn.ftz.f32 %f143, %f140, %f60, %f586; // begin inline asm { cvt.f32.f16 %f61, %rs10;} // end inline asm fma.rn.ftz.f32 %f144, %f140, %f61, %f585; // begin inline asm { cvt.f32.f16 %f62, %rs11;} // end inline asm fma.rn.ftz.f32 %f145, %f140, %f62, %f584; cvt.rn.f32.s16 %f146, %rs12; sub.ftz.f32 %f147, %f146, %f7; mul.ftz.f32 %f148, %f57, %f147; // begin inline asm { cvt.f32.f16 %f63, %rs13;} // end inline asm fma.rn.ftz.f32 %f149, %f148, %f63, %f141; // begin inline asm { cvt.f32.f16 %f64, %rs14;} // end inline asm fma.rn.ftz.f32 %f150, %f148, %f64, %f142; // begin inline asm { cvt.f32.f16 %f65, %rs15;} // end inline asm fma.rn.ftz.f32 %f151, %f148, %f65, %f143; // begin inline asm { cvt.f32.f16 %f66, %rs16;} // end inline asm fma.rn.ftz.f32 %f152, %f148, %f66, %f144; // begin inline asm { cvt.f32.f16 %f67, %rs17;} // end inline asm fma.rn.ftz.f32 %f153, %f148, %f67, %f145; cvt.rn.f32.s16 %f154, %rs18; sub.ftz.f32 %f155, %f154, %f7; mul.ftz.f32 %f156, %f57, %f155; // begin inline asm { cvt.f32.f16 %f68, %rs19;} // end inline asm fma.rn.ftz.f32 %f157, %f156, %f68, %f149; // begin inline asm { cvt.f32.f16 %f69, %rs20;} // end inline asm fma.rn.ftz.f32 %f158, %f156, %f69, %f150; // begin inline asm { cvt.f32.f16 %f70, %rs21;} // end inline asm fma.rn.ftz.f32 %f159, %f156, %f70, %f151; // begin inline asm { cvt.f32.f16 %f71, %rs22;} // end inline asm fma.rn.ftz.f32 %f160, %f156, %f71, %f152; // begin inline asm { cvt.f32.f16 %f72, %rs23;} // end inline asm fma.rn.ftz.f32 %f161, %f156, %f72, %f153; cvt.rn.f32.s16 %f162, %rs24; sub.ftz.f32 %f163, %f162, %f7; mul.ftz.f32 %f164, %f57, %f163; // begin inline asm { cvt.f32.f16 %f73, %rs25;} // end inline asm fma.rn.ftz.f32 %f165, %f164, %f73, %f157; // begin inline asm { cvt.f32.f16 %f74, %rs26;} // end inline asm fma.rn.ftz.f32 %f166, %f164, %f74, %f158; // begin inline asm { cvt.f32.f16 %f75, %rs27;} // end inline asm fma.rn.ftz.f32 %f167, %f164, %f75, %f159; // begin inline asm { cvt.f32.f16 %f76, %rs28;} // end inline asm fma.rn.ftz.f32 %f168, %f164, %f76, %f160; // begin inline asm { cvt.f32.f16 %f77, %rs29;} // end inline asm fma.rn.ftz.f32 %f169, %f164, %f77, %f161; cvt.rn.f32.s16 %f170, %rs30; sub.ftz.f32 %f171, %f170, %f7; mul.ftz.f32 %f172, %f57, %f171; // begin inline asm { cvt.f32.f16 %f78, %rs31;} // end inline asm fma.rn.ftz.f32 %f173, %f172, %f78, %f165; // begin inline asm { cvt.f32.f16 %f79, %rs33;} // end inline asm fma.rn.ftz.f32 %f174, %f172, %f79, %f166; // begin inline asm { cvt.f32.f16 %f80, %rs35;} // end inline asm fma.rn.ftz.f32 %f175, %f172, %f80, %f167; // begin inline asm { cvt.f32.f16 %f81, %rs37;} // end inline asm fma.rn.ftz.f32 %f176, %f172, %f81, %f168; // begin inline asm { cvt.f32.f16 %f82, %rs39;} // end inline asm fma.rn.ftz.f32 %f177, %f172, %f82, %f169; cvt.rn.f32.s16 %f178, %rs41; sub.ftz.f32 %f179, %f178, %f7; mul.ftz.f32 %f180, %f57, %f179; // begin inline asm { cvt.f32.f16 %f83, %rs42;} // end inline asm fma.rn.ftz.f32 %f181, %f180, %f83, %f173; // begin inline asm { cvt.f32.f16 %f84, %rs43;} // end inline asm fma.rn.ftz.f32 %f182, %f180, %f84, %f174; // begin inline asm { cvt.f32.f16 %f85, %rs44;} // end inline asm fma.rn.ftz.f32 %f183, %f180, %f85, %f175; // begin inline asm { cvt.f32.f16 %f86, %rs45;} // end inline asm fma.rn.ftz.f32 %f184, %f180, %f86, %f176; // begin inline asm { cvt.f32.f16 %f87, %rs46;} // end inline asm fma.rn.ftz.f32 %f185, %f180, %f87, %f177; cvt.rn.f32.s16 %f186, %rs47; sub.ftz.f32 %f187, %f186, %f7; mul.ftz.f32 %f188, %f57, %f187; mov.b32 {%rs124, %rs129}, %r68; // begin inline asm { cvt.f32.f16 %f88, %rs124;} // end inline asm fma.rn.ftz.f32 %f189, %f188, %f88, %f181; mov.b32 {%rs125, %rs130}, %r72; // begin inline asm { cvt.f32.f16 %f89, %rs125;} // end inline asm fma.rn.ftz.f32 %f190, %f188, %f89, %f182; mov.b32 {%rs126, %rs131}, %r79; // begin inline asm { cvt.f32.f16 %f90, %rs126;} // end inline asm fma.rn.ftz.f32 %f191, %f188, %f90, %f183; mov.b32 {%rs127, %rs132}, %r84; // begin inline asm { cvt.f32.f16 %f91, %rs127;} // end inline asm fma.rn.ftz.f32 %f192, %f188, %f91, %f184; mov.b32 {%rs128, %rs133}, %r88; // begin inline asm { cvt.f32.f16 %f92, %rs128;} // end inline asm fma.rn.ftz.f32 %f193, %f188, %f92, %f185; cvt.rn.f32.s16 %f194, %rs48; sub.ftz.f32 %f195, %f194, %f7; mul.ftz.f32 %f196, %f57, %f195; // begin inline asm { cvt.f32.f16 %f93, %rs129;} // end inline asm fma.rn.ftz.f32 %f197, %f196, %f93, %f189; // begin inline asm { cvt.f32.f16 %f94, %rs130;} // end inline asm fma.rn.ftz.f32 %f198, %f196, %f94, %f190; // begin inline asm { cvt.f32.f16 %f95, %rs131;} // end inline asm fma.rn.ftz.f32 %f199, %f196, %f95, %f191; // begin inline asm { cvt.f32.f16 %f96, %rs132;} // end inline asm fma.rn.ftz.f32 %f200, %f196, %f96, %f192; // begin inline asm { cvt.f32.f16 %f97, %rs133;} // end inline asm fma.rn.ftz.f32 %f201, %f196, %f97, %f193; ld.global.v4.u32 {%r100, %r101, %r102, %r103}, [%rd5+16]; ld.global.v4.u32 {%r108, %r109, %r110, %r111}, [%rd6+16]; ld.global.v4.u32 {%r116, %r117, %r118, %r119}, [%rd7]; ld.global.v4.u32 {%r124, %r125, %r126, %r127}, [%rd8]; ld.global.v4.u32 {%r132, %r133, %r134, %r135}, [%rd9]; cvt.rn.f32.s16 %f202, %rs50; sub.ftz.f32 %f203, %f202, %f7; mul.ftz.f32 %f204, %f57, %f203; mov.b32 {%rs134, %rs139}, %r100; // begin inline asm { cvt.f32.f16 %f98, %rs134;} // end inline asm fma.rn.ftz.f32 %f205, %f204, %f98, %f197; mov.b32 {%rs135, %rs140}, %r108; // begin inline asm { cvt.f32.f16 %f99, %rs135;} // end inline asm fma.rn.ftz.f32 %f206, %f204, %f99, %f198; mov.b32 {%rs136, %rs141}, %r116; // begin inline asm { cvt.f32.f16 %f100, %rs136;} // end inline asm fma.rn.ftz.f32 %f207, %f204, %f100, %f199; mov.b32 {%rs137, %rs142}, %r124; // begin inline asm { cvt.f32.f16 %f101, %rs137;} // end inline asm fma.rn.ftz.f32 %f208, %f204, %f101, %f200; mov.b32 {%rs138, %rs143}, %r132; // begin inline asm { cvt.f32.f16 %f102, %rs138;} // end inline asm fma.rn.ftz.f32 %f209, %f204, %f102, %f201; cvt.rn.f32.s16 %f210, %rs52; sub.ftz.f32 %f211, %f210, %f7; mul.ftz.f32 %f212, %f57, %f211; // begin inline asm { cvt.f32.f16 %f103, %rs139;} // end inline asm fma.rn.ftz.f32 %f213, %f212, %f103, %f205; // begin inline asm { cvt.f32.f16 %f104, %rs140;} // end inline asm fma.rn.ftz.f32 %f214, %f212, %f104, %f206; // begin inline asm { cvt.f32.f16 %f105, %rs141;} // end inline asm fma.rn.ftz.f32 %f215, %f212, %f105, %f207; // begin inline asm { cvt.f32.f16 %f106, %rs142;} // end inline asm fma.rn.ftz.f32 %f216, %f212, %f106, %f208; // begin inline asm { cvt.f32.f16 %f107, %rs143;} // end inline asm fma.rn.ftz.f32 %f217, %f212, %f107, %f209; cvt.rn.f32.s16 %f218, %rs54; sub.ftz.f32 %f219, %f218, %f7; mul.ftz.f32 %f220, %f57, %f219; mov.b32 {%rs144, %rs149}, %r101; // begin inline asm { cvt.f32.f16 %f108, %rs144;} // end inline asm fma.rn.ftz.f32 %f221, %f220, %f108, %f213; mov.b32 {%rs145, %rs150}, %r109; // begin inline asm { cvt.f32.f16 %f109, %rs145;} // end inline asm fma.rn.ftz.f32 %f222, %f220, %f109, %f214; mov.b32 {%rs146, %rs151}, %r117; // begin inline asm { cvt.f32.f16 %f110, %rs146;} // end inline asm fma.rn.ftz.f32 %f223, %f220, %f110, %f215; mov.b32 {%rs147, %rs152}, %r125; // begin inline asm { cvt.f32.f16 %f111, %rs147;} // end inline asm fma.rn.ftz.f32 %f224, %f220, %f111, %f216; mov.b32 {%rs148, %rs153}, %r133; // begin inline asm { cvt.f32.f16 %f112, %rs148;} // end inline asm fma.rn.ftz.f32 %f225, %f220, %f112, %f217; cvt.rn.f32.s16 %f226, %rs56; sub.ftz.f32 %f227, %f226, %f7; mul.ftz.f32 %f228, %f57, %f227; // begin inline asm { cvt.f32.f16 %f113, %rs149;} // end inline asm fma.rn.ftz.f32 %f229, %f228, %f113, %f221; // begin inline asm { cvt.f32.f16 %f114, %rs150;} // end inline asm fma.rn.ftz.f32 %f230, %f228, %f114, %f222; // begin inline asm { cvt.f32.f16 %f115, %rs151;} // end inline asm fma.rn.ftz.f32 %f231, %f228, %f115, %f223; // begin inline asm { cvt.f32.f16 %f116, %rs152;} // end inline asm fma.rn.ftz.f32 %f232, %f228, %f116, %f224; // begin inline asm { cvt.f32.f16 %f117, %rs153;} // end inline asm fma.rn.ftz.f32 %f233, %f228, %f117, %f225; cvt.rn.f32.s16 %f234, %rs58; sub.ftz.f32 %f235, %f234, %f7; mul.ftz.f32 %f236, %f57, %f235; mov.b32 {%rs154, %rs159}, %r102; // begin inline asm { cvt.f32.f16 %f118, %rs154;} // end inline asm fma.rn.ftz.f32 %f237, %f236, %f118, %f229; mov.b32 {%rs155, %rs160}, %r110; // begin inline asm { cvt.f32.f16 %f119, %rs155;} // end inline asm fma.rn.ftz.f32 %f238, %f236, %f119, %f230; mov.b32 {%rs156, %rs161}, %r118; // begin inline asm { cvt.f32.f16 %f120, %rs156;} // end inline asm fma.rn.ftz.f32 %f239, %f236, %f120, %f231; mov.b32 {%rs157, %rs162}, %r126; // begin inline asm { cvt.f32.f16 %f121, %rs157;} // end inline asm fma.rn.ftz.f32 %f240, %f236, %f121, %f232; mov.b32 {%rs158, %rs163}, %r134; // begin inline asm { cvt.f32.f16 %f122, %rs158;} // end inline asm fma.rn.ftz.f32 %f241, %f236, %f122, %f233; cvt.rn.f32.s16 %f242, %rs60; sub.ftz.f32 %f243, %f242, %f7; mul.ftz.f32 %f244, %f57, %f243; // begin inline asm { cvt.f32.f16 %f123, %rs159;} // end inline asm fma.rn.ftz.f32 %f245, %f244, %f123, %f237; // begin inline asm { cvt.f32.f16 %f124, %rs160;} // end inline asm fma.rn.ftz.f32 %f246, %f244, %f124, %f238; // begin inline asm { cvt.f32.f16 %f125, %rs161;} // end inline asm fma.rn.ftz.f32 %f247, %f244, %f125, %f239; // begin inline asm { cvt.f32.f16 %f126, %rs162;} // end inline asm fma.rn.ftz.f32 %f248, %f244, %f126, %f240; // begin inline asm { cvt.f32.f16 %f127, %rs163;} // end inline asm fma.rn.ftz.f32 %f249, %f244, %f127, %f241; cvt.rn.f32.s16 %f250, %rs62; sub.ftz.f32 %f251, %f250, %f7; mul.ftz.f32 %f252, %f57, %f251; mov.b32 {%rs164, %rs169}, %r103; // begin inline asm { cvt.f32.f16 %f128, %rs164;} // end inline asm fma.rn.ftz.f32 %f253, %f252, %f128, %f245; mov.b32 {%rs165, %rs170}, %r111; // begin inline asm { cvt.f32.f16 %f129, %rs165;} // end inline asm fma.rn.ftz.f32 %f254, %f252, %f129, %f246; mov.b32 {%rs166, %rs171}, %r119; // begin inline asm { cvt.f32.f16 %f130, %rs166;} // end inline asm fma.rn.ftz.f32 %f255, %f252, %f130, %f247; mov.b32 {%rs167, %rs172}, %r127; // begin inline asm { cvt.f32.f16 %f131, %rs167;} // end inline asm fma.rn.ftz.f32 %f256, %f252, %f131, %f248; mov.b32 {%rs168, %rs173}, %r135; // begin inline asm { cvt.f32.f16 %f132, %rs168;} // end inline asm fma.rn.ftz.f32 %f257, %f252, %f132, %f249; cvt.rn.f32.s16 %f258, %rs63; sub.ftz.f32 %f259, %f258, %f7; mul.ftz.f32 %f260, %f57, %f259; // begin inline asm { cvt.f32.f16 %f133, %rs169;} // end inline asm fma.rn.ftz.f32 %f588, %f260, %f133, %f253; // begin inline asm { cvt.f32.f16 %f134, %rs170;} // end inline asm fma.rn.ftz.f32 %f587, %f260, %f134, %f254; // begin inline asm { cvt.f32.f16 %f135, %rs171;} // end inline asm fma.rn.ftz.f32 %f586, %f260, %f135, %f255; // begin inline asm { cvt.f32.f16 %f136, %rs172;} // end inline asm fma.rn.ftz.f32 %f585, %f260, %f136, %f256; // begin inline asm { cvt.f32.f16 %f137, %rs173;} // end inline asm fma.rn.ftz.f32 %f584, %f260, %f137, %f257; bra.uni $L__BB0_8; $L__BB0_7: shl.b16 %rs255, %rs88, 4; cvt.s16.s8 %rs256, %rs255; shr.s16 %rs257, %rs256, 7; and.b16 %rs258, %rs257, -16; or.b16 %rs259, %rs258, %rs5; cvt.rn.f32.s16 %f341, %rs259; sub.ftz.f32 %f342, %f341, %f7; mul.ftz.f32 %f343, %f57, %f342; cvt.u16.u32 %rs174, %r65; // begin inline asm { cvt.f32.f16 %f261, %rs174;} // end inline asm fma.rn.ftz.f32 %f344, %f343, %f261, %f588; // begin inline asm { cvt.f32.f16 %f262, %rs8;} // end inline asm fma.rn.ftz.f32 %f345, %f343, %f262, %f587; // begin inline asm { cvt.f32.f16 %f263, %rs9;} // end inline asm fma.rn.ftz.f32 %f346, %f343, %f263, %f586; // begin inline asm { cvt.f32.f16 %f264, %rs10;} // end inline asm fma.rn.ftz.f32 %f347, %f343, %f264, %f585; // begin inline asm { cvt.f32.f16 %f265, %rs11;} // end inline asm fma.rn.ftz.f32 %f348, %f343, %f265, %f584; cvt.s16.s8 %rs260, %rs88; shr.s16 %rs261, %rs260, 7; and.b16 %rs262, %rs261, -16; or.b16 %rs263, %rs262, %rs12; cvt.rn.f32.s16 %f349, %rs263; sub.ftz.f32 %f350, %f349, %f7; mul.ftz.f32 %f351, %f57, %f350; // begin inline asm { cvt.f32.f16 %f266, %rs13;} // end inline asm fma.rn.ftz.f32 %f352, %f351, %f266, %f344; // begin inline asm { cvt.f32.f16 %f267, %rs14;} // end inline asm fma.rn.ftz.f32 %f353, %f351, %f267, %f345; // begin inline asm { cvt.f32.f16 %f268, %rs15;} // end inline asm fma.rn.ftz.f32 %f354, %f351, %f268, %f346; // begin inline asm { cvt.f32.f16 %f269, %rs16;} // end inline asm fma.rn.ftz.f32 %f355, %f351, %f269, %f347; // begin inline asm { cvt.f32.f16 %f270, %rs17;} // end inline asm fma.rn.ftz.f32 %f356, %f351, %f270, %f348; cvt.s16.s8 %rs265, %rs89; shr.s16 %rs266, %rs265, 7; and.b16 %rs267, %rs266, -16; or.b16 %rs268, %rs267, %rs18; cvt.rn.f32.s16 %f357, %rs268; sub.ftz.f32 %f358, %f357, %f7; mul.ftz.f32 %f359, %f57, %f358; // begin inline asm { cvt.f32.f16 %f271, %rs19;} // end inline asm fma.rn.ftz.f32 %f360, %f359, %f271, %f352; // begin inline asm { cvt.f32.f16 %f272, %rs20;} // end inline asm fma.rn.ftz.f32 %f361, %f359, %f272, %f353; // begin inline asm { cvt.f32.f16 %f273, %rs21;} // end inline asm fma.rn.ftz.f32 %f362, %f359, %f273, %f354; // begin inline asm { cvt.f32.f16 %f274, %rs22;} // end inline asm fma.rn.ftz.f32 %f363, %f359, %f274, %f355; // begin inline asm { cvt.f32.f16 %f275, %rs23;} // end inline asm fma.rn.ftz.f32 %f364, %f359, %f275, %f356; shr.s16 %rs269, %rs88, 15; and.b16 %rs270, %rs269, -16; or.b16 %rs271, %rs270, %rs24; cvt.rn.f32.s16 %f365, %rs271; sub.ftz.f32 %f366, %f365, %f7; mul.ftz.f32 %f367, %f57, %f366; // begin inline asm { cvt.f32.f16 %f276, %rs25;} // end inline asm fma.rn.ftz.f32 %f368, %f367, %f276, %f360; // begin inline asm { cvt.f32.f16 %f277, %rs26;} // end inline asm fma.rn.ftz.f32 %f369, %f367, %f277, %f361; // begin inline asm { cvt.f32.f16 %f278, %rs27;} // end inline asm fma.rn.ftz.f32 %f370, %f367, %f278, %f362; // begin inline asm { cvt.f32.f16 %f279, %rs28;} // end inline asm fma.rn.ftz.f32 %f371, %f367, %f279, %f363; // begin inline asm { cvt.f32.f16 %f280, %rs29;} // end inline asm fma.rn.ftz.f32 %f372, %f367, %f280, %f364; shl.b16 %rs273, %rs91, 4; cvt.s16.s8 %rs274, %rs273; shr.s16 %rs275, %rs274, 7; and.b16 %rs276, %rs275, -16; or.b16 %rs277, %rs276, %rs30; cvt.rn.f32.s16 %f373, %rs277; sub.ftz.f32 %f374, %f373, %f7; mul.ftz.f32 %f375, %f57, %f374; // begin inline asm { cvt.f32.f16 %f281, %rs31;} // end inline asm fma.rn.ftz.f32 %f376, %f375, %f281, %f368; // begin inline asm { cvt.f32.f16 %f282, %rs33;} // end inline asm fma.rn.ftz.f32 %f377, %f375, %f282, %f369; // begin inline asm { cvt.f32.f16 %f283, %rs35;} // end inline asm fma.rn.ftz.f32 %f378, %f375, %f283, %f370; // begin inline asm { cvt.f32.f16 %f284, %rs37;} // end inline asm fma.rn.ftz.f32 %f379, %f375, %f284, %f371; // begin inline asm { cvt.f32.f16 %f285, %rs39;} // end inline asm fma.rn.ftz.f32 %f380, %f375, %f285, %f372; shl.b16 %rs279, %rs92, 4; cvt.s16.s8 %rs280, %rs279; shr.s16 %rs281, %rs280, 7; and.b16 %rs282, %rs281, -16; or.b16 %rs283, %rs282, %rs41; cvt.rn.f32.s16 %f381, %rs283; sub.ftz.f32 %f382, %f381, %f7; mul.ftz.f32 %f383, %f57, %f382; // begin inline asm { cvt.f32.f16 %f286, %rs42;} // end inline asm fma.rn.ftz.f32 %f384, %f383, %f286, %f376; // begin inline asm { cvt.f32.f16 %f287, %rs43;} // end inline asm fma.rn.ftz.f32 %f385, %f383, %f287, %f377; // begin inline asm { cvt.f32.f16 %f288, %rs44;} // end inline asm fma.rn.ftz.f32 %f386, %f383, %f288, %f378; // begin inline asm { cvt.f32.f16 %f289, %rs45;} // end inline asm fma.rn.ftz.f32 %f387, %f383, %f289, %f379; // begin inline asm { cvt.f32.f16 %f290, %rs46;} // end inline asm fma.rn.ftz.f32 %f388, %f383, %f290, %f380; shl.b16 %rs285, %rs93, 4; cvt.s16.s8 %rs286, %rs285; shr.s16 %rs287, %rs286, 7; and.b16 %rs288, %rs287, -16; or.b16 %rs289, %rs288, %rs47; cvt.rn.f32.s16 %f389, %rs289; sub.ftz.f32 %f390, %f389, %f7; mul.ftz.f32 %f391, %f57, %f390; mov.b32 {%rs204, %rs209}, %r68; // begin inline asm { cvt.f32.f16 %f291, %rs204;} // end inline asm fma.rn.ftz.f32 %f392, %f391, %f291, %f384; mov.b32 {%rs205, %rs210}, %r72; // begin inline asm { cvt.f32.f16 %f292, %rs205;} // end inline asm fma.rn.ftz.f32 %f393, %f391, %f292, %f385; mov.b32 {%rs206, %rs211}, %r79; // begin inline asm { cvt.f32.f16 %f293, %rs206;} // end inline asm fma.rn.ftz.f32 %f394, %f391, %f293, %f386; mov.b32 {%rs207, %rs212}, %r84; // begin inline asm { cvt.f32.f16 %f294, %rs207;} // end inline asm fma.rn.ftz.f32 %f395, %f391, %f294, %f387; mov.b32 {%rs208, %rs213}, %r88; // begin inline asm { cvt.f32.f16 %f295, %rs208;} // end inline asm fma.rn.ftz.f32 %f396, %f391, %f295, %f388; shl.b16 %rs290, %rs48, 4; cvt.s16.s8 %rs291, %rs290; shr.s16 %rs292, %rs291, 7; and.b16 %rs293, %rs292, -16; or.b16 %rs294, %rs293, %rs48; cvt.rn.f32.s16 %f397, %rs294; sub.ftz.f32 %f398, %f397, %f7; mul.ftz.f32 %f399, %f57, %f398; // begin inline asm { cvt.f32.f16 %f296, %rs209;} // end inline asm fma.rn.ftz.f32 %f400, %f399, %f296, %f392; // begin inline asm { cvt.f32.f16 %f297, %rs210;} // end inline asm fma.rn.ftz.f32 %f401, %f399, %f297, %f393; // begin inline asm { cvt.f32.f16 %f298, %rs211;} // end inline asm fma.rn.ftz.f32 %f402, %f399, %f298, %f394; // begin inline asm { cvt.f32.f16 %f299, %rs212;} // end inline asm fma.rn.ftz.f32 %f403, %f399, %f299, %f395; // begin inline asm { cvt.f32.f16 %f300, %rs213;} // end inline asm fma.rn.ftz.f32 %f404, %f399, %f300, %f396; ld.global.v4.u32 {%r143, %r144, %r145, %r146}, [%rd5+16]; ld.global.v4.u32 {%r151, %r152, %r153, %r154}, [%rd6+16]; ld.global.v4.u32 {%r159, %r160, %r161, %r162}, [%rd7]; ld.global.v4.u32 {%r167, %r168, %r169, %r170}, [%rd8]; ld.global.v4.u32 {%r175, %r176, %r177, %r178}, [%rd9]; shl.b16 %rs295, %rs49, 4; cvt.s16.s8 %rs296, %rs295; shr.s16 %rs297, %rs296, 7; and.b16 %rs298, %rs297, -16; or.b16 %rs299, %rs298, %rs50; cvt.rn.f32.s16 %f405, %rs299; sub.ftz.f32 %f406, %f405, %f7; mul.ftz.f32 %f407, %f57, %f406; mov.b32 {%rs214, %rs219}, %r143; // begin inline asm { cvt.f32.f16 %f301, %rs214;} // end inline asm fma.rn.ftz.f32 %f408, %f407, %f301, %f400; mov.b32 {%rs215, %rs220}, %r151; // begin inline asm { cvt.f32.f16 %f302, %rs215;} // end inline asm fma.rn.ftz.f32 %f409, %f407, %f302, %f401; mov.b32 {%rs216, %rs221}, %r159; // begin inline asm { cvt.f32.f16 %f303, %rs216;} // end inline asm fma.rn.ftz.f32 %f410, %f407, %f303, %f402; mov.b32 {%rs217, %rs222}, %r167; // begin inline asm { cvt.f32.f16 %f304, %rs217;} // end inline asm fma.rn.ftz.f32 %f411, %f407, %f304, %f403; mov.b32 {%rs218, %rs223}, %r175; // begin inline asm { cvt.f32.f16 %f305, %rs218;} // end inline asm fma.rn.ftz.f32 %f412, %f407, %f305, %f404; shl.b16 %rs300, %rs51, 4; cvt.s16.s8 %rs301, %rs300; shr.s16 %rs302, %rs301, 7; and.b16 %rs303, %rs302, -16; or.b16 %rs304, %rs303, %rs52; cvt.rn.f32.s16 %f413, %rs304; sub.ftz.f32 %f414, %f413, %f7; mul.ftz.f32 %f415, %f57, %f414; // begin inline asm { cvt.f32.f16 %f306, %rs219;} // end inline asm fma.rn.ftz.f32 %f416, %f415, %f306, %f408; // begin inline asm { cvt.f32.f16 %f307, %rs220;} // end inline asm fma.rn.ftz.f32 %f417, %f415, %f307, %f409; // begin inline asm { cvt.f32.f16 %f308, %rs221;} // end inline asm fma.rn.ftz.f32 %f418, %f415, %f308, %f410; // begin inline asm { cvt.f32.f16 %f309, %rs222;} // end inline asm fma.rn.ftz.f32 %f419, %f415, %f309, %f411; // begin inline asm { cvt.f32.f16 %f310, %rs223;} // end inline asm fma.rn.ftz.f32 %f420, %f415, %f310, %f412; shl.b16 %rs305, %rs53, 4; cvt.s16.s8 %rs306, %rs305; shr.s16 %rs307, %rs306, 7; and.b16 %rs308, %rs307, -16; or.b16 %rs309, %rs308, %rs54; cvt.rn.f32.s16 %f421, %rs309; sub.ftz.f32 %f422, %f421, %f7; mul.ftz.f32 %f423, %f57, %f422; mov.b32 {%rs224, %rs229}, %r144; // begin inline asm { cvt.f32.f16 %f311, %rs224;} // end inline asm fma.rn.ftz.f32 %f424, %f423, %f311, %f416; mov.b32 {%rs225, %rs230}, %r152; // begin inline asm { cvt.f32.f16 %f312, %rs225;} // end inline asm fma.rn.ftz.f32 %f425, %f423, %f312, %f417; mov.b32 {%rs226, %rs231}, %r160; // begin inline asm { cvt.f32.f16 %f313, %rs226;} // end inline asm fma.rn.ftz.f32 %f426, %f423, %f313, %f418; mov.b32 {%rs227, %rs232}, %r168; // begin inline asm { cvt.f32.f16 %f314, %rs227;} // end inline asm fma.rn.ftz.f32 %f427, %f423, %f314, %f419; mov.b32 {%rs228, %rs233}, %r176; // begin inline asm { cvt.f32.f16 %f315, %rs228;} // end inline asm fma.rn.ftz.f32 %f428, %f423, %f315, %f420; shl.b16 %rs310, %rs55, 4; cvt.s16.s8 %rs311, %rs310; shr.s16 %rs312, %rs311, 7; and.b16 %rs313, %rs312, -16; or.b16 %rs314, %rs313, %rs56; cvt.rn.f32.s16 %f429, %rs314; sub.ftz.f32 %f430, %f429, %f7; mul.ftz.f32 %f431, %f57, %f430; // begin inline asm { cvt.f32.f16 %f316, %rs229;} // end inline asm fma.rn.ftz.f32 %f432, %f431, %f316, %f424; // begin inline asm { cvt.f32.f16 %f317, %rs230;} // end inline asm fma.rn.ftz.f32 %f433, %f431, %f317, %f425; // begin inline asm { cvt.f32.f16 %f318, %rs231;} // end inline asm fma.rn.ftz.f32 %f434, %f431, %f318, %f426; // begin inline asm { cvt.f32.f16 %f319, %rs232;} // end inline asm fma.rn.ftz.f32 %f435, %f431, %f319, %f427; // begin inline asm { cvt.f32.f16 %f320, %rs233;} // end inline asm fma.rn.ftz.f32 %f436, %f431, %f320, %f428; shl.b16 %rs315, %rs57, 4; cvt.s16.s8 %rs316, %rs315; shr.s16 %rs317, %rs316, 7; and.b16 %rs318, %rs317, -16; or.b16 %rs319, %rs318, %rs58; cvt.rn.f32.s16 %f437, %rs319; sub.ftz.f32 %f438, %f437, %f7; mul.ftz.f32 %f439, %f57, %f438; mov.b32 {%rs234, %rs239}, %r145; // begin inline asm { cvt.f32.f16 %f321, %rs234;} // end inline asm fma.rn.ftz.f32 %f440, %f439, %f321, %f432; mov.b32 {%rs235, %rs240}, %r153; // begin inline asm { cvt.f32.f16 %f322, %rs235;} // end inline asm fma.rn.ftz.f32 %f441, %f439, %f322, %f433; mov.b32 {%rs236, %rs241}, %r161; // begin inline asm { cvt.f32.f16 %f323, %rs236;} // end inline asm fma.rn.ftz.f32 %f442, %f439, %f323, %f434; mov.b32 {%rs237, %rs242}, %r169; // begin inline asm { cvt.f32.f16 %f324, %rs237;} // end inline asm fma.rn.ftz.f32 %f443, %f439, %f324, %f435; mov.b32 {%rs238, %rs243}, %r177; // begin inline asm { cvt.f32.f16 %f325, %rs238;} // end inline asm fma.rn.ftz.f32 %f444, %f439, %f325, %f436; shl.b16 %rs320, %rs59, 4; cvt.s16.s8 %rs321, %rs320; shr.s16 %rs322, %rs321, 7; and.b16 %rs323, %rs322, -16; or.b16 %rs324, %rs323, %rs60; cvt.rn.f32.s16 %f445, %rs324; sub.ftz.f32 %f446, %f445, %f7; mul.ftz.f32 %f447, %f57, %f446; // begin inline asm { cvt.f32.f16 %f326, %rs239;} // end inline asm fma.rn.ftz.f32 %f448, %f447, %f326, %f440; // begin inline asm { cvt.f32.f16 %f327, %rs240;} // end inline asm fma.rn.ftz.f32 %f449, %f447, %f327, %f441; // begin inline asm { cvt.f32.f16 %f328, %rs241;} // end inline asm fma.rn.ftz.f32 %f450, %f447, %f328, %f442; // begin inline asm { cvt.f32.f16 %f329, %rs242;} // end inline asm fma.rn.ftz.f32 %f451, %f447, %f329, %f443; // begin inline asm { cvt.f32.f16 %f330, %rs243;} // end inline asm fma.rn.ftz.f32 %f452, %f447, %f330, %f444; shl.b16 %rs325, %rs61, 4; cvt.s16.s8 %rs326, %rs325; shr.s16 %rs327, %rs326, 7; and.b16 %rs328, %rs327, -16; or.b16 %rs329, %rs328, %rs62; cvt.rn.f32.s16 %f453, %rs329; sub.ftz.f32 %f454, %f453, %f7; mul.ftz.f32 %f455, %f57, %f454; mov.b32 {%rs244, %rs249}, %r146; // begin inline asm { cvt.f32.f16 %f331, %rs244;} // end inline asm fma.rn.ftz.f32 %f456, %f455, %f331, %f448; mov.b32 {%rs245, %rs250}, %r154; // begin inline asm { cvt.f32.f16 %f332, %rs245;} // end inline asm fma.rn.ftz.f32 %f457, %f455, %f332, %f449; mov.b32 {%rs246, %rs251}, %r162; // begin inline asm { cvt.f32.f16 %f333, %rs246;} // end inline asm fma.rn.ftz.f32 %f458, %f455, %f333, %f450; mov.b32 {%rs247, %rs252}, %r170; // begin inline asm { cvt.f32.f16 %f334, %rs247;} // end inline asm fma.rn.ftz.f32 %f459, %f455, %f334, %f451; mov.b32 {%rs248, %rs253}, %r178; // begin inline asm { cvt.f32.f16 %f335, %rs248;} // end inline asm fma.rn.ftz.f32 %f460, %f455, %f335, %f452; shl.b16 %rs330, %rs63, 4; cvt.s16.s8 %rs331, %rs330; shr.s16 %rs332, %rs331, 7; and.b16 %rs333, %rs332, -16; or.b16 %rs334, %rs333, %rs63; cvt.rn.f32.s16 %f461, %rs334; sub.ftz.f32 %f462, %f461, %f7; mul.ftz.f32 %f463, %f57, %f462; // begin inline asm { cvt.f32.f16 %f336, %rs249;} // end inline asm fma.rn.ftz.f32 %f588, %f463, %f336, %f456; // begin inline asm { cvt.f32.f16 %f337, %rs250;} // end inline asm fma.rn.ftz.f32 %f587, %f463, %f337, %f457; // begin inline asm { cvt.f32.f16 %f338, %rs251;} // end inline asm fma.rn.ftz.f32 %f586, %f463, %f338, %f458; // begin inline asm { cvt.f32.f16 %f339, %rs252;} // end inline asm fma.rn.ftz.f32 %f585, %f463, %f339, %f459; // begin inline asm { cvt.f32.f16 %f340, %rs253;} // end inline asm fma.rn.ftz.f32 %f584, %f463, %f340, %f460; $L__BB0_8: add.s32 %r290, %r290, 4; shl.b32 %r183, %r290, 5; add.s32 %r289, %r183, %r49; shl.b32 %r288, %r289, 1; setp.lt.u32 %p7, %r288, %r46; @%p7 bra $L__BB0_2; $L__BB0_9: mov.u32 %r283, %tid.y; mov.u32 %r282, %tid.x; shl.b32 %r281, %r283, 5; add.s32 %r280, %r281, %r282; shl.b32 %r185, %r280, 2; mov.u32 %r186, _ZZ9gemv_int4ILi4ELi64ELi5EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage; add.s32 %r187, %r186, %r185; setp.lt.u32 %p8, %r280, 32; @%p8 bra $L__BB0_11; add.s32 %r275, %r187, -112; st.shared.f32 [%r275], %f588; $L__BB0_11: mov.u32 %r287, %tid.y; mov.u32 %r286, %tid.x; shl.b32 %r285, %r287, 5; add.s32 %r284, %r285, %r286; setp.gt.u32 %p9, %r284, 31; bar.sync 0; mad.lo.s32 %r39, %r284, 12, %r186; @%p9 bra $L__BB0_13; mov.u32 %r202, 16; ld.shared.f32 %f479, [%r39+16]; add.ftz.f32 %f480, %f588, %f479; ld.shared.f32 %f481, [%r39+20]; add.ftz.f32 %f482, %f480, %f481; ld.shared.f32 %f483, [%r39+24]; add.ftz.f32 %f466, %f482, %f483; mov.u32 %r190, 1; mov.u32 %r203, 31; mov.u32 %r204, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f466, %r190, %r203, %r204; @p add.f32 r0, r0, %f466; mov.f32 %f464, r0;} // end inline asm mov.u32 %r193, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f464, %r193, %r203, %r204; @p add.f32 r0, r0, %f464; mov.f32 %f467, r0;} // end inline asm mov.u32 %r196, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f467, %r196, %r203, %r204; @p add.f32 r0, r0, %f467; mov.f32 %f470, r0;} // end inline asm mov.u32 %r199, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f470, %r199, %r203, %r204; @p add.f32 r0, r0, %f470; mov.f32 %f473, r0;} // end inline asm // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f473, %r202, %r203, %r204; @p add.f32 r0, r0, %f473; mov.f32 %f588, r0;} // end inline asm $L__BB0_13: @%p8 bra $L__BB0_15; add.s32 %r276, %r187, -112; st.shared.f32 [%r276+640], %f587; $L__BB0_15: bar.sync 0; @%p9 bra $L__BB0_17; ld.shared.f32 %f499, [%r39+656]; add.ftz.f32 %f500, %f587, %f499; ld.shared.f32 %f501, [%r39+660]; add.ftz.f32 %f502, %f500, %f501; ld.shared.f32 %f503, [%r39+664]; add.ftz.f32 %f486, %f502, %f503; mov.u32 %r206, 1; mov.u32 %r219, 31; mov.u32 %r220, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f486, %r206, %r219, %r220; @p add.f32 r0, r0, %f486; mov.f32 %f484, r0;} // end inline asm mov.u32 %r209, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f484, %r209, %r219, %r220; @p add.f32 r0, r0, %f484; mov.f32 %f487, r0;} // end inline asm mov.u32 %r212, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f487, %r212, %r219, %r220; @p add.f32 r0, r0, %f487; mov.f32 %f490, r0;} // end inline asm mov.u32 %r215, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f490, %r215, %r219, %r220; @p add.f32 r0, r0, %f490; mov.f32 %f493, r0;} // end inline asm mov.u32 %r218, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f493, %r218, %r219, %r220; @p add.f32 r0, r0, %f493; mov.f32 %f587, r0;} // end inline asm $L__BB0_17: @%p8 bra $L__BB0_19; add.s32 %r277, %r187, -112; st.shared.f32 [%r277+1280], %f586; $L__BB0_19: bar.sync 0; @%p9 bra $L__BB0_21; ld.shared.f32 %f519, [%r39+1296]; add.ftz.f32 %f520, %f586, %f519; ld.shared.f32 %f521, [%r39+1300]; add.ftz.f32 %f522, %f520, %f521; ld.shared.f32 %f523, [%r39+1304]; add.ftz.f32 %f506, %f522, %f523; mov.u32 %r222, 1; mov.u32 %r235, 31; mov.u32 %r236, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f506, %r222, %r235, %r236; @p add.f32 r0, r0, %f506; mov.f32 %f504, r0;} // end inline asm mov.u32 %r225, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f504, %r225, %r235, %r236; @p add.f32 r0, r0, %f504; mov.f32 %f507, r0;} // end inline asm mov.u32 %r228, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f507, %r228, %r235, %r236; @p add.f32 r0, r0, %f507; mov.f32 %f510, r0;} // end inline asm mov.u32 %r231, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f510, %r231, %r235, %r236; @p add.f32 r0, r0, %f510; mov.f32 %f513, r0;} // end inline asm mov.u32 %r234, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f513, %r234, %r235, %r236; @p add.f32 r0, r0, %f513; mov.f32 %f586, r0;} // end inline asm $L__BB0_21: @%p8 bra $L__BB0_23; add.s32 %r278, %r187, -112; st.shared.f32 [%r278+1920], %f585; $L__BB0_23: bar.sync 0; @%p9 bra $L__BB0_25; ld.shared.f32 %f539, [%r39+1936]; add.ftz.f32 %f540, %f585, %f539; ld.shared.f32 %f541, [%r39+1940]; add.ftz.f32 %f542, %f540, %f541; ld.shared.f32 %f543, [%r39+1944]; add.ftz.f32 %f526, %f542, %f543; mov.u32 %r238, 1; mov.u32 %r251, 31; mov.u32 %r252, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f526, %r238, %r251, %r252; @p add.f32 r0, r0, %f526; mov.f32 %f524, r0;} // end inline asm mov.u32 %r241, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f524, %r241, %r251, %r252; @p add.f32 r0, r0, %f524; mov.f32 %f527, r0;} // end inline asm mov.u32 %r244, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f527, %r244, %r251, %r252; @p add.f32 r0, r0, %f527; mov.f32 %f530, r0;} // end inline asm mov.u32 %r247, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f530, %r247, %r251, %r252; @p add.f32 r0, r0, %f530; mov.f32 %f533, r0;} // end inline asm mov.u32 %r250, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f533, %r250, %r251, %r252; @p add.f32 r0, r0, %f533; mov.f32 %f585, r0;} // end inline asm $L__BB0_25: @%p8 bra $L__BB0_27; add.s32 %r279, %r187, -112; st.shared.f32 [%r279+2560], %f584; $L__BB0_27: bar.sync 0; @%p9 bra $L__BB0_29; ld.shared.f32 %f559, [%r39+2576]; add.ftz.f32 %f560, %f584, %f559; ld.shared.f32 %f561, [%r39+2580]; add.ftz.f32 %f562, %f560, %f561; ld.shared.f32 %f563, [%r39+2584]; add.ftz.f32 %f546, %f562, %f563; mov.u32 %r254, 1; mov.u32 %r267, 31; mov.u32 %r268, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f546, %r254, %r267, %r268; @p add.f32 r0, r0, %f546; mov.f32 %f544, r0;} // end inline asm mov.u32 %r257, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f544, %r257, %r267, %r268; @p add.f32 r0, r0, %f544; mov.f32 %f547, r0;} // end inline asm mov.u32 %r260, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f547, %r260, %r267, %r268; @p add.f32 r0, r0, %f547; mov.f32 %f550, r0;} // end inline asm mov.u32 %r263, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f550, %r263, %r267, %r268; @p add.f32 r0, r0, %f550; mov.f32 %f553, r0;} // end inline asm mov.u32 %r266, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f553, %r266, %r267, %r268; @p add.f32 r0, r0, %f553; mov.f32 %f584, r0;} // end inline asm $L__BB0_29: mov.u32 %r269, %tid.y; or.b32 %r271, %r49, %r269; setp.ne.s32 %p18, %r271, 0; @%p18 bra $L__BB0_33; ld.param.u64 %rd58, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+8]; ld.param.u64 %rd57, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0]; cvta.to.global.u64 %rd39, %rd57; setp.eq.s64 %p19, %rd58, 0; mul.ftz.f32 %f38, %f50, %f588; mov.u32 %r272, %ctaid.x; cvt.s64.s32 %rd11, %r272; mul.wide.s32 %rd40, %r272, 2; add.s64 %rd12, %rd39, %rd40; mul.ftz.f32 %f39, %f50, %f587; add.s32 %r273, %r45, %r272; cvt.s64.s32 %rd13, %r45; mul.wide.s32 %rd41, %r45, 2; add.s64 %rd14, %rd12, %rd41; mul.ftz.f32 %f40, %f50, %f586; add.s32 %r274, %r273, %r45; cvt.s64.s32 %rd15, %r274; mul.wide.s32 %rd42, %r274, 2; add.s64 %rd17, %rd39, %rd42; mul.ftz.f32 %f41, %f50, %f585; mul.ftz.f32 %f42, %f50, %f584; @%p19 bra $L__BB0_32; ld.param.u64 %rd59, [_Z27dequant_gemv_group64_batch523DequantGemvKernelParams_param_0+8]; cvta.to.global.u64 %rd43, %rd59; shl.b64 %rd44, %rd11, 1; add.s64 %rd45, %rd43, %rd44; ld.global.u16 %rs335, [%rd45]; // begin inline asm { cvt.f32.f16 %f564, %rs335;} // end inline asm fma.rn.ftz.f32 %f565, %f51, %f564, %f38; // begin inline asm { cvt.rn.f16.f32 %rs336, %f565;} // end inline asm st.global.u16 [%rd12], %rs336; shl.b64 %rd46, %rd13, 1; add.s64 %rd47, %rd45, %rd46; ld.global.u16 %rs337, [%rd47]; // begin inline asm { cvt.f32.f16 %f566, %rs337;} // end inline asm fma.rn.ftz.f32 %f567, %f51, %f566, %f39; // begin inline asm { cvt.rn.f16.f32 %rs338, %f567;} // end inline asm st.global.u16 [%rd14], %rs338; shl.b64 %rd48, %rd15, 1; add.s64 %rd49, %rd43, %rd48; ld.global.u16 %rs339, [%rd49]; // begin inline asm { cvt.f32.f16 %f568, %rs339;} // end inline asm fma.rn.ftz.f32 %f569, %f51, %f568, %f40; // begin inline asm { cvt.rn.f16.f32 %rs340, %f569;} // end inline asm st.global.u16 [%rd17], %rs340; add.s64 %rd50, %rd49, %rd46; ld.global.u16 %rs341, [%rd50]; // begin inline asm { cvt.f32.f16 %f570, %rs341;} // end inline asm fma.rn.ftz.f32 %f571, %f51, %f570, %f41; // begin inline asm { cvt.rn.f16.f32 %rs342, %f571;} // end inline asm add.s64 %rd51, %rd17, %rd46; st.global.u16 [%rd51], %rs342; add.s64 %rd52, %rd50, %rd46; ld.global.u16 %rs343, [%rd52]; // begin inline asm { cvt.f32.f16 %f572, %rs343;} // end inline asm fma.rn.ftz.f32 %f573, %f51, %f572, %f42; // begin inline asm { cvt.rn.f16.f32 %rs344, %f573;} // end inline asm add.s64 %rd53, %rd51, %rd46; st.global.u16 [%rd53], %rs344; bra.uni $L__BB0_33; $L__BB0_32: // begin inline asm { cvt.rn.f16.f32 %rs345, %f38;} // end inline asm st.global.u16 [%rd12], %rs345; // begin inline asm { cvt.rn.f16.f32 %rs346, %f39;} // end inline asm st.global.u16 [%rd14], %rs346; // begin inline asm { cvt.rn.f16.f32 %rs347, %f40;} // end inline asm st.global.u16 [%rd17], %rs347; // begin inline asm { cvt.rn.f16.f32 %rs348, %f41;} // end inline asm shl.b64 %rd54, %rd13, 1; add.s64 %rd55, %rd17, %rd54; st.global.u16 [%rd55], %rs348; // begin inline asm { cvt.rn.f16.f32 %rs349, %f42;} // end inline asm add.s64 %rd56, %rd55, %rd54; st.global.u16 [%rd56], %rs349; $L__BB0_33: ret; } // .globl _ZN3cub11EmptyKernelIvEEvv .visible .entry _ZN3cub11EmptyKernelIvEEvv() { ret; }