85 // Based on NVVM 7.0.1 // .version 7.8 .target sm_60 .address_size 64 // .globl _Z27dequant_gemv_group64_batch223DequantGemvKernelParams // _ZZ9gemv_int4ILi4ELi64ELi2EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage has been demoted .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust6system6detail10sequential3seqE[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_1E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_2E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_3E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_4E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_5E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_6E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_7E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_8E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders2_9E[1]; .global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_e50bdfa76thrust12placeholders3_10E[1]; .visible .entry _Z27dequant_gemv_group64_batch223DequantGemvKernelParams( .param .align 8 .b8 _Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0[80] ) { .reg .pred %p<14>; .reg .b16 %rs<226>; .reg .f32 %f<302>; .reg .b32 %r<157>; .reg .b64 %rd<38>; // demoted variable .shared .align 16 .b8 _ZZ9gemv_int4ILi4ELi64ELi2EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage[1280]; ld.param.v2.u32 {%r28, %r29}, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+48]; ld.param.v2.u32 {%r30, %r31}, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+56]; ld.param.v2.f32 {%f23, %f24}, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+64]; ld.param.v4.u8 {%rs59, %rs60, %rs61, %rs62}, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+72]; ld.param.u64 %rd18, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+40]; ld.param.u64 %rd17, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+32]; ld.param.u64 %rd16, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+24]; ld.param.u64 %rd15, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+16]; mov.u32 %r1, %ctaid.x; mov.u32 %r156, %tid.y; shl.b32 %r32, %r156, 5; mov.u32 %r3, %tid.x; add.s32 %r155, %r32, %r3; shl.b32 %r154, %r155, 1; setp.ge.u32 %p1, %r154, %r30; mov.f32 %f296, 0f00000000; mov.f32 %f297, %f296; @%p1 bra $L__BB0_9; mul.lo.s32 %r8, %r30, %r1; shr.u32 %r9, %r3, 2; mul.lo.s32 %r10, %r31, %r1; shl.b16 %rs2, %rs59, 3; cvta.to.global.u64 %rd2, %rd18; cvta.to.global.u64 %rd3, %rd17; cvta.to.global.u64 %rd4, %rd16; cvta.to.global.u64 %rd5, %rd15; $L__BB0_2: add.s32 %r33, %r154, %r8; mul.wide.u32 %rd19, %r33, 4; add.s64 %rd20, %rd4, %rd19; ld.global.v2.u32 {%r34, %r35}, [%rd20]; shl.b32 %r36, %r156, 3; add.s32 %r16, %r36, %r9; add.s32 %r17, %r16, %r10; mul.wide.s32 %rd21, %r17, 2; add.s64 %rd22, %rd2, %rd21; ld.global.u16 %rs67, [%rd22]; // begin inline asm { cvt.f32.f16 %f27, %rs67;} // end inline asm setp.eq.s64 %p2, %rd17, 0; mov.u16 %rs225, %rs2; @%p2 bra $L__BB0_4; shr.u32 %r37, %r17, 31; add.s32 %r38, %r17, %r37; shr.s32 %r39, %r38, 1; cvt.s64.s32 %rd23, %r39; add.s64 %rd24, %rd3, %rd23; ld.global.u8 %r40, [%rd24]; shl.b32 %r41, %r16, 2; and.b32 %r42, %r41, 4; shr.u32 %r43, %r40, %r42; cvt.u16.u32 %rs68, %r43; and.b16 %rs225, %rs68, 15; $L__BB0_4: shl.b32 %r18, %r155, 4; setp.ge.s32 %p3, %r18, %r28; @%p3 bra $L__BB0_8; setp.eq.s16 %p4, %rs59, 0; shr.u16 %rs70, %rs225, 3; and.b16 %rs71, %rs70, 1; setp.eq.b16 %p5, %rs71, 1; and.pred %p6, %p4, %p5; selp.b16 %rs72, -16, 0, %p6; or.b16 %rs73, %rs72, %rs225; cvt.s16.s8 %rs74, %rs73; cvt.rn.f32.s16 %f4, %rs74; mul.wide.s32 %rd25, %r18, 2; add.s64 %rd6, %rd5, %rd25; ld.global.v4.u32 {%r44, %r45, %r46, %r47}, [%rd6]; mul.wide.s32 %rd26, %r28, 2; add.s64 %rd7, %rd6, %rd26; ld.global.v4.u32 {%r52, %r53, %r54, %r55}, [%rd7]; cvt.u16.u32 %rs5, %r34; and.b16 %rs6, %rs5, 15; mov.b32 {%rs7, %rs11}, %r44; mov.b32 {%rs8, %rs12}, %r52; shr.u32 %r60, %r34, 4; cvt.u16.u32 %rs9, %r60; and.b16 %rs10, %rs9, 15; shr.u32 %r61, %r34, 8; cvt.u16.u32 %rs13, %r61; and.b16 %rs14, %rs13, 15; mov.b32 {%rs15, %rs19}, %r45; mov.b32 {%rs16, %rs20}, %r53; shr.u32 %r62, %r34, 12; cvt.u16.u32 %rs17, %r62; and.b16 %rs18, %rs17, 15; shr.u32 %r63, %r34, 16; cvt.u16.u32 %rs21, %r63; and.b16 %rs22, %rs21, 15; mov.b32 {%rs23, %rs27}, %r46; mov.b32 {%rs24, %rs28}, %r54; shr.u32 %r64, %r34, 20; cvt.u16.u32 %rs25, %r64; and.b16 %rs26, %rs25, 15; shr.u32 %r65, %r34, 24; cvt.u16.u32 %rs29, %r65; and.b16 %rs30, %rs29, 15; mov.b32 {%rs31, %rs34}, %r47; mov.b32 {%rs32, %rs35}, %r55; shr.u32 %r66, %r34, 28; cvt.u16.u32 %rs33, %r66; cvt.u16.u32 %rs36, %r35; and.b16 %rs37, %rs36, 15; shr.u32 %r67, %r35, 4; cvt.u16.u32 %rs38, %r67; and.b16 %rs39, %rs38, 15; shr.u32 %r68, %r35, 8; cvt.u16.u32 %rs40, %r68; and.b16 %rs41, %rs40, 15; shr.u32 %r69, %r35, 12; cvt.u16.u32 %rs42, %r69; and.b16 %rs43, %rs42, 15; shr.u32 %r70, %r35, 16; cvt.u16.u32 %rs44, %r70; and.b16 %rs45, %rs44, 15; shr.u32 %r71, %r35, 20; cvt.u16.u32 %rs46, %r71; and.b16 %rs47, %rs46, 15; shr.u32 %r72, %r35, 24; cvt.u16.u32 %rs48, %r72; and.b16 %rs49, %rs48, 15; shr.u32 %r73, %r35, 28; cvt.u16.u32 %rs50, %r73; @%p4 bra $L__BB0_7; cvt.rn.f32.s16 %f60, %rs6; sub.ftz.f32 %f61, %f60, %f4; mul.ftz.f32 %f62, %f27, %f61; // begin inline asm { cvt.f32.f16 %f28, %rs7;} // end inline asm fma.rn.ftz.f32 %f63, %f62, %f28, %f297; // begin inline asm { cvt.f32.f16 %f29, %rs8;} // end inline asm fma.rn.ftz.f32 %f64, %f62, %f29, %f296; cvt.rn.f32.s16 %f65, %rs10; sub.ftz.f32 %f66, %f65, %f4; mul.ftz.f32 %f67, %f27, %f66; // begin inline asm { cvt.f32.f16 %f30, %rs11;} // end inline asm fma.rn.ftz.f32 %f68, %f67, %f30, %f63; // begin inline asm { cvt.f32.f16 %f31, %rs12;} // end inline asm fma.rn.ftz.f32 %f69, %f67, %f31, %f64; cvt.rn.f32.s16 %f70, %rs14; sub.ftz.f32 %f71, %f70, %f4; mul.ftz.f32 %f72, %f27, %f71; // begin inline asm { cvt.f32.f16 %f32, %rs15;} // end inline asm fma.rn.ftz.f32 %f73, %f72, %f32, %f68; // begin inline asm { cvt.f32.f16 %f33, %rs16;} // end inline asm fma.rn.ftz.f32 %f74, %f72, %f33, %f69; cvt.rn.f32.s16 %f75, %rs18; sub.ftz.f32 %f76, %f75, %f4; mul.ftz.f32 %f77, %f27, %f76; // begin inline asm { cvt.f32.f16 %f34, %rs19;} // end inline asm fma.rn.ftz.f32 %f78, %f77, %f34, %f73; // begin inline asm { cvt.f32.f16 %f35, %rs20;} // end inline asm fma.rn.ftz.f32 %f79, %f77, %f35, %f74; cvt.rn.f32.s16 %f80, %rs22; sub.ftz.f32 %f81, %f80, %f4; mul.ftz.f32 %f82, %f27, %f81; // begin inline asm { cvt.f32.f16 %f36, %rs23;} // end inline asm fma.rn.ftz.f32 %f83, %f82, %f36, %f78; // begin inline asm { cvt.f32.f16 %f37, %rs24;} // end inline asm fma.rn.ftz.f32 %f84, %f82, %f37, %f79; cvt.rn.f32.s16 %f85, %rs26; sub.ftz.f32 %f86, %f85, %f4; mul.ftz.f32 %f87, %f27, %f86; // begin inline asm { cvt.f32.f16 %f38, %rs27;} // end inline asm fma.rn.ftz.f32 %f88, %f87, %f38, %f83; // begin inline asm { cvt.f32.f16 %f39, %rs28;} // end inline asm fma.rn.ftz.f32 %f89, %f87, %f39, %f84; cvt.rn.f32.s16 %f90, %rs30; sub.ftz.f32 %f91, %f90, %f4; mul.ftz.f32 %f92, %f27, %f91; // begin inline asm { cvt.f32.f16 %f40, %rs31;} // end inline asm fma.rn.ftz.f32 %f93, %f92, %f40, %f88; // begin inline asm { cvt.f32.f16 %f41, %rs32;} // end inline asm fma.rn.ftz.f32 %f94, %f92, %f41, %f89; cvt.rn.f32.s16 %f95, %rs33; sub.ftz.f32 %f96, %f95, %f4; mul.ftz.f32 %f97, %f27, %f96; // begin inline asm { cvt.f32.f16 %f42, %rs34;} // end inline asm fma.rn.ftz.f32 %f98, %f97, %f42, %f93; // begin inline asm { cvt.f32.f16 %f43, %rs35;} // end inline asm fma.rn.ftz.f32 %f99, %f97, %f43, %f94; ld.global.v4.u32 {%r74, %r75, %r76, %r77}, [%rd6+16]; ld.global.v4.u32 {%r82, %r83, %r84, %r85}, [%rd7+16]; cvt.rn.f32.s16 %f100, %rs37; sub.ftz.f32 %f101, %f100, %f4; mul.ftz.f32 %f102, %f27, %f101; mov.b32 {%rs91, %rs93}, %r74; // begin inline asm { cvt.f32.f16 %f44, %rs91;} // end inline asm fma.rn.ftz.f32 %f103, %f102, %f44, %f98; mov.b32 {%rs92, %rs94}, %r82; // begin inline asm { cvt.f32.f16 %f45, %rs92;} // end inline asm fma.rn.ftz.f32 %f104, %f102, %f45, %f99; cvt.rn.f32.s16 %f105, %rs39; sub.ftz.f32 %f106, %f105, %f4; mul.ftz.f32 %f107, %f27, %f106; // begin inline asm { cvt.f32.f16 %f46, %rs93;} // end inline asm fma.rn.ftz.f32 %f108, %f107, %f46, %f103; // begin inline asm { cvt.f32.f16 %f47, %rs94;} // end inline asm fma.rn.ftz.f32 %f109, %f107, %f47, %f104; cvt.rn.f32.s16 %f110, %rs41; sub.ftz.f32 %f111, %f110, %f4; mul.ftz.f32 %f112, %f27, %f111; mov.b32 {%rs95, %rs97}, %r75; // begin inline asm { cvt.f32.f16 %f48, %rs95;} // end inline asm fma.rn.ftz.f32 %f113, %f112, %f48, %f108; mov.b32 {%rs96, %rs98}, %r83; // begin inline asm { cvt.f32.f16 %f49, %rs96;} // end inline asm fma.rn.ftz.f32 %f114, %f112, %f49, %f109; cvt.rn.f32.s16 %f115, %rs43; sub.ftz.f32 %f116, %f115, %f4; mul.ftz.f32 %f117, %f27, %f116; // begin inline asm { cvt.f32.f16 %f50, %rs97;} // end inline asm fma.rn.ftz.f32 %f118, %f117, %f50, %f113; // begin inline asm { cvt.f32.f16 %f51, %rs98;} // end inline asm fma.rn.ftz.f32 %f119, %f117, %f51, %f114; cvt.rn.f32.s16 %f120, %rs45; sub.ftz.f32 %f121, %f120, %f4; mul.ftz.f32 %f122, %f27, %f121; mov.b32 {%rs99, %rs101}, %r76; // begin inline asm { cvt.f32.f16 %f52, %rs99;} // end inline asm fma.rn.ftz.f32 %f123, %f122, %f52, %f118; mov.b32 {%rs100, %rs102}, %r84; // begin inline asm { cvt.f32.f16 %f53, %rs100;} // end inline asm fma.rn.ftz.f32 %f124, %f122, %f53, %f119; cvt.rn.f32.s16 %f125, %rs47; sub.ftz.f32 %f126, %f125, %f4; mul.ftz.f32 %f127, %f27, %f126; // begin inline asm { cvt.f32.f16 %f54, %rs101;} // end inline asm fma.rn.ftz.f32 %f128, %f127, %f54, %f123; // begin inline asm { cvt.f32.f16 %f55, %rs102;} // end inline asm fma.rn.ftz.f32 %f129, %f127, %f55, %f124; cvt.rn.f32.s16 %f130, %rs49; sub.ftz.f32 %f131, %f130, %f4; mul.ftz.f32 %f132, %f27, %f131; mov.b32 {%rs103, %rs105}, %r77; // begin inline asm { cvt.f32.f16 %f56, %rs103;} // end inline asm fma.rn.ftz.f32 %f133, %f132, %f56, %f128; mov.b32 {%rs104, %rs106}, %r85; // begin inline asm { cvt.f32.f16 %f57, %rs104;} // end inline asm fma.rn.ftz.f32 %f134, %f132, %f57, %f129; cvt.rn.f32.s16 %f135, %rs50; sub.ftz.f32 %f136, %f135, %f4; mul.ftz.f32 %f137, %f27, %f136; // begin inline asm { cvt.f32.f16 %f58, %rs105;} // end inline asm fma.rn.ftz.f32 %f297, %f137, %f58, %f133; // begin inline asm { cvt.f32.f16 %f59, %rs106;} // end inline asm fma.rn.ftz.f32 %f296, %f137, %f59, %f134; bra.uni $L__BB0_8; $L__BB0_7: shl.b16 %rs139, %rs5, 4; cvt.s16.s8 %rs140, %rs139; shr.s16 %rs141, %rs140, 7; and.b16 %rs142, %rs141, -16; or.b16 %rs143, %rs142, %rs6; cvt.rn.f32.s16 %f170, %rs143; sub.ftz.f32 %f171, %f170, %f4; mul.ftz.f32 %f172, %f27, %f171; // begin inline asm { cvt.f32.f16 %f138, %rs7;} // end inline asm fma.rn.ftz.f32 %f173, %f172, %f138, %f297; // begin inline asm { cvt.f32.f16 %f139, %rs8;} // end inline asm fma.rn.ftz.f32 %f174, %f172, %f139, %f296; shl.b16 %rs144, %rs9, 4; cvt.s16.s8 %rs145, %rs144; shr.s16 %rs146, %rs145, 7; and.b16 %rs147, %rs146, -16; or.b16 %rs148, %rs147, %rs10; cvt.rn.f32.s16 %f175, %rs148; sub.ftz.f32 %f176, %f175, %f4; mul.ftz.f32 %f177, %f27, %f176; // begin inline asm { cvt.f32.f16 %f140, %rs11;} // end inline asm fma.rn.ftz.f32 %f178, %f177, %f140, %f173; // begin inline asm { cvt.f32.f16 %f141, %rs12;} // end inline asm fma.rn.ftz.f32 %f179, %f177, %f141, %f174; shl.b16 %rs149, %rs13, 4; cvt.s16.s8 %rs150, %rs149; shr.s16 %rs151, %rs150, 7; and.b16 %rs152, %rs151, -16; or.b16 %rs153, %rs152, %rs14; cvt.rn.f32.s16 %f180, %rs153; sub.ftz.f32 %f181, %f180, %f4; mul.ftz.f32 %f182, %f27, %f181; // begin inline asm { cvt.f32.f16 %f142, %rs15;} // end inline asm fma.rn.ftz.f32 %f183, %f182, %f142, %f178; // begin inline asm { cvt.f32.f16 %f143, %rs16;} // end inline asm fma.rn.ftz.f32 %f184, %f182, %f143, %f179; shl.b16 %rs154, %rs17, 4; cvt.s16.s8 %rs155, %rs154; shr.s16 %rs156, %rs155, 7; and.b16 %rs157, %rs156, -16; or.b16 %rs158, %rs157, %rs18; cvt.rn.f32.s16 %f185, %rs158; sub.ftz.f32 %f186, %f185, %f4; mul.ftz.f32 %f187, %f27, %f186; // begin inline asm { cvt.f32.f16 %f144, %rs19;} // end inline asm fma.rn.ftz.f32 %f188, %f187, %f144, %f183; // begin inline asm { cvt.f32.f16 %f145, %rs20;} // end inline asm fma.rn.ftz.f32 %f189, %f187, %f145, %f184; shl.b16 %rs159, %rs21, 4; cvt.s16.s8 %rs160, %rs159; shr.s16 %rs161, %rs160, 7; and.b16 %rs162, %rs161, -16; or.b16 %rs163, %rs162, %rs22; cvt.rn.f32.s16 %f190, %rs163; sub.ftz.f32 %f191, %f190, %f4; mul.ftz.f32 %f192, %f27, %f191; // begin inline asm { cvt.f32.f16 %f146, %rs23;} // end inline asm fma.rn.ftz.f32 %f193, %f192, %f146, %f188; // begin inline asm { cvt.f32.f16 %f147, %rs24;} // end inline asm fma.rn.ftz.f32 %f194, %f192, %f147, %f189; shl.b16 %rs164, %rs25, 4; cvt.s16.s8 %rs165, %rs164; shr.s16 %rs166, %rs165, 7; and.b16 %rs167, %rs166, -16; or.b16 %rs168, %rs167, %rs26; cvt.rn.f32.s16 %f195, %rs168; sub.ftz.f32 %f196, %f195, %f4; mul.ftz.f32 %f197, %f27, %f196; // begin inline asm { cvt.f32.f16 %f148, %rs27;} // end inline asm fma.rn.ftz.f32 %f198, %f197, %f148, %f193; // begin inline asm { cvt.f32.f16 %f149, %rs28;} // end inline asm fma.rn.ftz.f32 %f199, %f197, %f149, %f194; shl.b16 %rs169, %rs29, 4; cvt.s16.s8 %rs170, %rs169; shr.s16 %rs171, %rs170, 7; and.b16 %rs172, %rs171, -16; or.b16 %rs173, %rs172, %rs30; cvt.rn.f32.s16 %f200, %rs173; sub.ftz.f32 %f201, %f200, %f4; mul.ftz.f32 %f202, %f27, %f201; // begin inline asm { cvt.f32.f16 %f150, %rs31;} // end inline asm fma.rn.ftz.f32 %f203, %f202, %f150, %f198; // begin inline asm { cvt.f32.f16 %f151, %rs32;} // end inline asm fma.rn.ftz.f32 %f204, %f202, %f151, %f199; shl.b16 %rs174, %rs33, 4; cvt.s16.s8 %rs175, %rs174; shr.s16 %rs176, %rs175, 7; and.b16 %rs177, %rs176, -16; or.b16 %rs178, %rs177, %rs33; cvt.rn.f32.s16 %f205, %rs178; sub.ftz.f32 %f206, %f205, %f4; mul.ftz.f32 %f207, %f27, %f206; // begin inline asm { cvt.f32.f16 %f152, %rs34;} // end inline asm fma.rn.ftz.f32 %f208, %f207, %f152, %f203; // begin inline asm { cvt.f32.f16 %f153, %rs35;} // end inline asm fma.rn.ftz.f32 %f209, %f207, %f153, %f204; ld.global.v4.u32 {%r90, %r91, %r92, %r93}, [%rd6+16]; ld.global.v4.u32 {%r98, %r99, %r100, %r101}, [%rd7+16]; shl.b16 %rs179, %rs36, 4; cvt.s16.s8 %rs180, %rs179; shr.s16 %rs181, %rs180, 7; and.b16 %rs182, %rs181, -16; or.b16 %rs183, %rs182, %rs37; cvt.rn.f32.s16 %f210, %rs183; sub.ftz.f32 %f211, %f210, %f4; mul.ftz.f32 %f212, %f27, %f211; mov.b32 {%rs123, %rs125}, %r90; // begin inline asm { cvt.f32.f16 %f154, %rs123;} // end inline asm fma.rn.ftz.f32 %f213, %f212, %f154, %f208; mov.b32 {%rs124, %rs126}, %r98; // begin inline asm { cvt.f32.f16 %f155, %rs124;} // end inline asm fma.rn.ftz.f32 %f214, %f212, %f155, %f209; shl.b16 %rs184, %rs38, 4; cvt.s16.s8 %rs185, %rs184; shr.s16 %rs186, %rs185, 7; and.b16 %rs187, %rs186, -16; or.b16 %rs188, %rs187, %rs39; cvt.rn.f32.s16 %f215, %rs188; sub.ftz.f32 %f216, %f215, %f4; mul.ftz.f32 %f217, %f27, %f216; // begin inline asm { cvt.f32.f16 %f156, %rs125;} // end inline asm fma.rn.ftz.f32 %f218, %f217, %f156, %f213; // begin inline asm { cvt.f32.f16 %f157, %rs126;} // end inline asm fma.rn.ftz.f32 %f219, %f217, %f157, %f214; shl.b16 %rs189, %rs40, 4; cvt.s16.s8 %rs190, %rs189; shr.s16 %rs191, %rs190, 7; and.b16 %rs192, %rs191, -16; or.b16 %rs193, %rs192, %rs41; cvt.rn.f32.s16 %f220, %rs193; sub.ftz.f32 %f221, %f220, %f4; mul.ftz.f32 %f222, %f27, %f221; mov.b32 {%rs127, %rs129}, %r91; // begin inline asm { cvt.f32.f16 %f158, %rs127;} // end inline asm fma.rn.ftz.f32 %f223, %f222, %f158, %f218; mov.b32 {%rs128, %rs130}, %r99; // begin inline asm { cvt.f32.f16 %f159, %rs128;} // end inline asm fma.rn.ftz.f32 %f224, %f222, %f159, %f219; shl.b16 %rs194, %rs42, 4; cvt.s16.s8 %rs195, %rs194; shr.s16 %rs196, %rs195, 7; and.b16 %rs197, %rs196, -16; or.b16 %rs198, %rs197, %rs43; cvt.rn.f32.s16 %f225, %rs198; sub.ftz.f32 %f226, %f225, %f4; mul.ftz.f32 %f227, %f27, %f226; // begin inline asm { cvt.f32.f16 %f160, %rs129;} // end inline asm fma.rn.ftz.f32 %f228, %f227, %f160, %f223; // begin inline asm { cvt.f32.f16 %f161, %rs130;} // end inline asm fma.rn.ftz.f32 %f229, %f227, %f161, %f224; shl.b16 %rs199, %rs44, 4; cvt.s16.s8 %rs200, %rs199; shr.s16 %rs201, %rs200, 7; and.b16 %rs202, %rs201, -16; or.b16 %rs203, %rs202, %rs45; cvt.rn.f32.s16 %f230, %rs203; sub.ftz.f32 %f231, %f230, %f4; mul.ftz.f32 %f232, %f27, %f231; mov.b32 {%rs131, %rs133}, %r92; // begin inline asm { cvt.f32.f16 %f162, %rs131;} // end inline asm fma.rn.ftz.f32 %f233, %f232, %f162, %f228; mov.b32 {%rs132, %rs134}, %r100; // begin inline asm { cvt.f32.f16 %f163, %rs132;} // end inline asm fma.rn.ftz.f32 %f234, %f232, %f163, %f229; shl.b16 %rs204, %rs46, 4; cvt.s16.s8 %rs205, %rs204; shr.s16 %rs206, %rs205, 7; and.b16 %rs207, %rs206, -16; or.b16 %rs208, %rs207, %rs47; cvt.rn.f32.s16 %f235, %rs208; sub.ftz.f32 %f236, %f235, %f4; mul.ftz.f32 %f237, %f27, %f236; // begin inline asm { cvt.f32.f16 %f164, %rs133;} // end inline asm fma.rn.ftz.f32 %f238, %f237, %f164, %f233; // begin inline asm { cvt.f32.f16 %f165, %rs134;} // end inline asm fma.rn.ftz.f32 %f239, %f237, %f165, %f234; shl.b16 %rs209, %rs48, 4; cvt.s16.s8 %rs210, %rs209; shr.s16 %rs211, %rs210, 7; and.b16 %rs212, %rs211, -16; or.b16 %rs213, %rs212, %rs49; cvt.rn.f32.s16 %f240, %rs213; sub.ftz.f32 %f241, %f240, %f4; mul.ftz.f32 %f242, %f27, %f241; mov.b32 {%rs135, %rs137}, %r93; // begin inline asm { cvt.f32.f16 %f166, %rs135;} // end inline asm fma.rn.ftz.f32 %f243, %f242, %f166, %f238; mov.b32 {%rs136, %rs138}, %r101; // begin inline asm { cvt.f32.f16 %f167, %rs136;} // end inline asm fma.rn.ftz.f32 %f244, %f242, %f167, %f239; shl.b16 %rs214, %rs50, 4; cvt.s16.s8 %rs215, %rs214; shr.s16 %rs216, %rs215, 7; and.b16 %rs217, %rs216, -16; or.b16 %rs218, %rs217, %rs50; cvt.rn.f32.s16 %f245, %rs218; sub.ftz.f32 %f246, %f245, %f4; mul.ftz.f32 %f247, %f27, %f246; // begin inline asm { cvt.f32.f16 %f168, %rs137;} // end inline asm fma.rn.ftz.f32 %f297, %f247, %f168, %f243; // begin inline asm { cvt.f32.f16 %f169, %rs138;} // end inline asm fma.rn.ftz.f32 %f296, %f247, %f169, %f244; $L__BB0_8: add.s32 %r156, %r156, 4; shl.b32 %r106, %r156, 5; add.s32 %r155, %r106, %r3; shl.b32 %r154, %r155, 1; setp.lt.u32 %p7, %r154, %r30; @%p7 bra $L__BB0_2; $L__BB0_9: mov.u32 %r150, %tid.y; shl.b32 %r149, %r150, 5; add.s32 %r148, %r149, %r3; shl.b32 %r107, %r148, 2; mov.u32 %r108, _ZZ9gemv_int4ILi4ELi64ELi2EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage; add.s32 %r109, %r108, %r107; setp.lt.u32 %p8, %r148, 32; @%p8 bra $L__BB0_11; add.s32 %r144, %r109, -112; st.shared.f32 [%r144], %f297; $L__BB0_11: mov.u32 %r153, %tid.y; shl.b32 %r152, %r153, 5; add.s32 %r151, %r152, %r3; setp.gt.u32 %p9, %r151, 31; bar.sync 0; mad.lo.s32 %r23, %r151, 12, %r108; @%p9 bra $L__BB0_13; mov.u32 %r124, 16; ld.shared.f32 %f263, [%r23+16]; add.ftz.f32 %f264, %f297, %f263; ld.shared.f32 %f265, [%r23+20]; add.ftz.f32 %f266, %f264, %f265; ld.shared.f32 %f267, [%r23+24]; add.ftz.f32 %f250, %f266, %f267; mov.u32 %r112, 1; mov.u32 %r125, 31; mov.u32 %r126, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f250, %r112, %r125, %r126; @p add.f32 r0, r0, %f250; mov.f32 %f248, r0;} // end inline asm mov.u32 %r115, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f248, %r115, %r125, %r126; @p add.f32 r0, r0, %f248; mov.f32 %f251, r0;} // end inline asm mov.u32 %r118, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f251, %r118, %r125, %r126; @p add.f32 r0, r0, %f251; mov.f32 %f254, r0;} // end inline asm mov.u32 %r121, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f254, %r121, %r125, %r126; @p add.f32 r0, r0, %f254; mov.f32 %f257, r0;} // end inline asm // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f257, %r124, %r125, %r126; @p add.f32 r0, r0, %f257; mov.f32 %f297, r0;} // end inline asm $L__BB0_13: @%p8 bra $L__BB0_15; add.s32 %r145, %r109, -112; st.shared.f32 [%r145+640], %f296; $L__BB0_15: bar.sync 0; @%p9 bra $L__BB0_17; ld.shared.f32 %f283, [%r23+656]; add.ftz.f32 %f284, %f296, %f283; ld.shared.f32 %f285, [%r23+660]; add.ftz.f32 %f286, %f284, %f285; ld.shared.f32 %f287, [%r23+664]; add.ftz.f32 %f270, %f286, %f287; mov.u32 %r128, 1; mov.u32 %r141, 31; mov.u32 %r142, -1; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f270, %r128, %r141, %r142; @p add.f32 r0, r0, %f270; mov.f32 %f268, r0;} // end inline asm mov.u32 %r131, 2; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f268, %r131, %r141, %r142; @p add.f32 r0, r0, %f268; mov.f32 %f271, r0;} // end inline asm mov.u32 %r134, 4; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f271, %r134, %r141, %r142; @p add.f32 r0, r0, %f271; mov.f32 %f274, r0;} // end inline asm mov.u32 %r137, 8; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f274, %r137, %r141, %r142; @p add.f32 r0, r0, %f274; mov.f32 %f277, r0;} // end inline asm mov.u32 %r140, 16; // begin inline asm { .reg .f32 r0; .reg .pred p; shfl.sync.down.b32 r0|p, %f277, %r140, %r141, %r142; @p add.f32 r0, r0, %f277; mov.f32 %f296, r0;} // end inline asm $L__BB0_17: mov.u32 %r146, %tid.y; or.b32 %r143, %r3, %r146; setp.ne.s32 %p12, %r143, 0; @%p12 bra $L__BB0_21; ld.param.u64 %rd36, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0]; ld.param.u64 %rd35, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+8]; mov.u32 %r147, %ctaid.x; setp.eq.s64 %p13, %rd35, 0; mul.ftz.f32 %f17, %f23, %f297; cvt.s64.s32 %rd9, %r147; cvta.to.global.u64 %rd27, %rd36; mul.wide.s32 %rd28, %r147, 2; add.s64 %rd10, %rd27, %rd28; mul.ftz.f32 %f18, %f23, %f296; cvt.s64.s32 %rd11, %r29; mul.wide.s32 %rd29, %r29, 2; add.s64 %rd12, %rd10, %rd29; @%p13 bra $L__BB0_20; ld.param.u64 %rd37, [_Z27dequant_gemv_group64_batch223DequantGemvKernelParams_param_0+8]; cvta.to.global.u64 %rd30, %rd37; shl.b64 %rd31, %rd9, 1; add.s64 %rd32, %rd30, %rd31; ld.global.u16 %rs219, [%rd32]; // begin inline asm { cvt.f32.f16 %f288, %rs219;} // end inline asm fma.rn.ftz.f32 %f289, %f24, %f288, %f17; // begin inline asm { cvt.rn.f16.f32 %rs220, %f289;} // end inline asm st.global.u16 [%rd10], %rs220; shl.b64 %rd33, %rd11, 1; add.s64 %rd34, %rd32, %rd33; ld.global.u16 %rs221, [%rd34]; // begin inline asm { cvt.f32.f16 %f290, %rs221;} // end inline asm fma.rn.ftz.f32 %f291, %f24, %f290, %f18; // begin inline asm { cvt.rn.f16.f32 %rs222, %f291;} // end inline asm st.global.u16 [%rd12], %rs222; bra.uni $L__BB0_21; $L__BB0_20: // begin inline asm { cvt.rn.f16.f32 %rs223, %f17;} // end inline asm st.global.u16 [%rd10], %rs223; // begin inline asm { cvt.rn.f16.f32 %rs224, %f18;} // end inline asm st.global.u16 [%rd12], %rs224; $L__BB0_21: ret; } // .globl _ZN3cub11EmptyKernelIvEEvv .visible .entry _ZN3cub11EmptyKernelIvEEvv() { ret; }