| // | |
| // Generated by NVIDIA NVVM Compiler | |
| // | |
| // Compiler Build ID: CL-33961263 | |
| // Cuda compilation tools, release 12.4, V12.4.99 | |
| // Based on NVVM 7.0.1 | |
| // | |
| .version 8.4 | |
| .target sm_52 | |
| .address_size 64 | |
| // .globl qpola_kernel_fp32 | |
| // _ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction has been demoted | |
| // _ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active has been demoted | |
| // _ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction has been demoted | |
| // _ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active has been demoted | |
| // _ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction has been demoted | |
| // _ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active has been demoted | |
| // _ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction has been demoted | |
| // _ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active has been demoted | |
| // _ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction has been demoted | |
| // _ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active has been demoted | |
| // _ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction has been demoted | |
| // _ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active has been demoted | |
| .visible .entry qpola_kernel_fp32( | |
| .param .u64 qpola_kernel_fp32_param_0, | |
| .param .u64 qpola_kernel_fp32_param_1, | |
| .param .f32 qpola_kernel_fp32_param_2, | |
| .param .f32 qpola_kernel_fp32_param_3, | |
| .param .u32 qpola_kernel_fp32_param_4 | |
| ) | |
| { | |
| .reg .pred %p<30>; | |
| .reg .f32 %f<135>; | |
| .reg .b32 %r<61>; | |
| .reg .b64 %rd<9>; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction[32]; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active[32]; | |
| ld.param.u64 %rd3, [qpola_kernel_fp32_param_0]; | |
| ld.param.u64 %rd2, [qpola_kernel_fp32_param_1]; | |
| ld.param.f32 %f14, [qpola_kernel_fp32_param_2]; | |
| ld.param.f32 %f15, [qpola_kernel_fp32_param_3]; | |
| ld.param.u32 %r6, [qpola_kernel_fp32_param_4]; | |
| cvta.to.global.u64 %rd4, %rd3; | |
| mov.u32 %r7, %ntid.x; | |
| mov.u32 %r8, %ctaid.x; | |
| mov.u32 %r1, %tid.x; | |
| mad.lo.s32 %r2, %r8, %r7, %r1; | |
| setp.lt.s32 %p1, %r2, %r6; | |
| setp.ge.s32 %p3, %r2, %r6; | |
| mul.wide.s32 %rd5, %r2, 4; | |
| add.s64 %rd1, %rd4, %rd5; | |
| mov.f32 %f132, 0f00000000; | |
| mov.f32 %f133, %f132; | |
| @%p3 bra $L__BB0_2; | |
| cvta.to.global.u64 %rd6, %rd2; | |
| add.s64 %rd8, %rd6, %rd5; | |
| ld.global.f32 %f132, [%rd8]; | |
| ld.global.f32 %f133, [%rd1]; | |
| $L__BB0_2: | |
| selp.f32 %f18, 0f3F800000, 0f00000000, %p1; | |
| abs.f32 %f19, %f132; | |
| setp.lt.f32 %p4, %f19, 0f7F800000; | |
| selp.f32 %f5, %f132, 0f00000000, %p4; | |
| setp.gt.f32 %p5, %f5, 0f00000000; | |
| setp.lt.f32 %p6, %f5, 0f00000000; | |
| selp.f32 %f20, 0fBF800000, 0f00000000, %p6; | |
| selp.f32 %f6, 0f3F800000, %f20, %p5; | |
| abs.f32 %f21, %f5; | |
| mov.b32 %r9, %f18; | |
| mov.u32 %r10, 2; | |
| mov.u32 %r11, 31; | |
| mov.u32 %r12, 16; | |
| mov.u32 %r13, -1; | |
| shfl.sync.down.b32 %r14|%p7, %r9, %r12, %r11, %r13; | |
| mov.b32 %f22, %r14; | |
| add.f32 %f23, %f18, %f22; | |
| mov.b32 %r15, %f6; | |
| shfl.sync.down.b32 %r16|%p8, %r15, %r12, %r11, %r13; | |
| mov.b32 %f24, %r16; | |
| add.f32 %f25, %f6, %f24; | |
| mov.b32 %r17, %f21; | |
| shfl.sync.down.b32 %r18|%p9, %r17, %r12, %r11, %r13; | |
| mov.b32 %f26, %r18; | |
| add.f32 %f27, %f21, %f26; | |
| mov.b32 %r19, %f23; | |
| mov.u32 %r20, 8; | |
| shfl.sync.down.b32 %r21|%p10, %r19, %r20, %r11, %r13; | |
| mov.b32 %f28, %r21; | |
| add.f32 %f29, %f23, %f28; | |
| mov.b32 %r22, %f25; | |
| shfl.sync.down.b32 %r23|%p11, %r22, %r20, %r11, %r13; | |
| mov.b32 %f30, %r23; | |
| add.f32 %f31, %f25, %f30; | |
| mov.b32 %r24, %f27; | |
| shfl.sync.down.b32 %r25|%p12, %r24, %r20, %r11, %r13; | |
| mov.b32 %f32, %r25; | |
| add.f32 %f33, %f27, %f32; | |
| mov.b32 %r26, %f29; | |
| mov.u32 %r27, 4; | |
| shfl.sync.down.b32 %r28|%p13, %r26, %r27, %r11, %r13; | |
| mov.b32 %f34, %r28; | |
| add.f32 %f35, %f29, %f34; | |
| mov.b32 %r29, %f31; | |
| shfl.sync.down.b32 %r30|%p14, %r29, %r27, %r11, %r13; | |
| mov.b32 %f36, %r30; | |
| add.f32 %f37, %f31, %f36; | |
| mov.b32 %r31, %f33; | |
| shfl.sync.down.b32 %r32|%p15, %r31, %r27, %r11, %r13; | |
| mov.b32 %f38, %r32; | |
| add.f32 %f39, %f33, %f38; | |
| mov.b32 %r33, %f35; | |
| shfl.sync.down.b32 %r34|%p16, %r33, %r10, %r11, %r13; | |
| mov.b32 %f40, %r34; | |
| add.f32 %f41, %f35, %f40; | |
| mov.b32 %r35, %f37; | |
| shfl.sync.down.b32 %r36|%p17, %r35, %r10, %r11, %r13; | |
| mov.b32 %f42, %r36; | |
| add.f32 %f43, %f37, %f42; | |
| mov.b32 %r37, %f39; | |
| shfl.sync.down.b32 %r38|%p18, %r37, %r10, %r11, %r13; | |
| mov.b32 %f44, %r38; | |
| add.f32 %f45, %f39, %f44; | |
| mov.b32 %r39, %f41; | |
| mov.u32 %r40, 1; | |
| shfl.sync.down.b32 %r41|%p19, %r39, %r40, %r11, %r13; | |
| mov.b32 %f46, %r41; | |
| add.f32 %f47, %f41, %f46; | |
| mov.b32 %r42, %f43; | |
| shfl.sync.down.b32 %r43|%p20, %r42, %r40, %r11, %r13; | |
| mov.b32 %f48, %r43; | |
| add.f32 %f49, %f43, %f48; | |
| mov.b32 %r44, %f45; | |
| shfl.sync.down.b32 %r45|%p21, %r44, %r40, %r11, %r13; | |
| mov.b32 %f50, %r45; | |
| add.f32 %f51, %f45, %f50; | |
| mov.b32 %r46, %f47; | |
| mov.u32 %r47, 0; | |
| shfl.sync.idx.b32 %r3|%p22, %r46, %r47, %r11, %r13; | |
| mov.b32 %r48, %f49; | |
| shfl.sync.idx.b32 %r4|%p23, %r48, %r47, %r11, %r13; | |
| mov.b32 %r49, %f51; | |
| shfl.sync.idx.b32 %r5|%p2, %r49, %r47, %r11, %r13; | |
| and.b32 %r50, %r1, 31; | |
| setp.ne.s32 %p24, %r50, 0; | |
| @%p24 bra $L__BB0_4; | |
| shr.u32 %r51, %r1, 3; | |
| and.b32 %r52, %r51, 536870908; | |
| mov.u32 %r53, _ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction; | |
| add.s32 %r54, %r53, %r52; | |
| st.shared.u32 [%r54], %r4; | |
| mov.u32 %r55, _ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active; | |
| add.s32 %r56, %r55, %r52; | |
| st.shared.u32 [%r56], %r3; | |
| $L__BB0_4: | |
| bar.sync 0; | |
| @%p3 bra $L__BB0_9; | |
| abs.f32 %f52, %f133; | |
| setp.lt.f32 %p26, %f52, 0f7F800000; | |
| selp.f32 %f7, %f133, 0f00000000, %p26; | |
| mov.f32 %f53, 0f00000000; | |
| mov.b32 %f54, %r3; | |
| mov.f32 %f55, 0f3F800000; | |
| max.f32 %f56, %f54, %f55; | |
| ld.shared.f32 %f57, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction]; | |
| ld.shared.f32 %f58, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active]; | |
| ld.shared.f32 %f59, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+4]; | |
| ld.shared.f32 %f60, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+4]; | |
| ld.shared.f32 %f61, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+8]; | |
| ld.shared.f32 %f62, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+8]; | |
| ld.shared.f32 %f63, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+12]; | |
| ld.shared.f32 %f64, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+12]; | |
| ld.shared.f32 %f65, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+16]; | |
| ld.shared.f32 %f66, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+16]; | |
| ld.shared.f32 %f67, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+20]; | |
| ld.shared.f32 %f68, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+20]; | |
| ld.shared.f32 %f69, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+24]; | |
| ld.shared.f32 %f70, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+24]; | |
| ld.shared.f32 %f71, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE17s_block_direction+28]; | |
| ld.shared.f32 %f72, [_ZZ17qpola_kernel_implIfEvPT_PKS0_ffiE14s_block_active+28]; | |
| mov.b32 %f73, %r4; | |
| div.rn.f32 %f74, %f73, %f56; | |
| mul.f32 %f75, %f6, %f74; | |
| sub.f32 %f76, %f55, %f75; | |
| max.f32 %f77, %f53, %f76; | |
| add.f32 %f78, %f58, 0f00000000; | |
| add.f32 %f79, %f78, %f60; | |
| add.f32 %f80, %f79, %f62; | |
| add.f32 %f81, %f80, %f64; | |
| add.f32 %f82, %f81, %f66; | |
| add.f32 %f83, %f82, %f68; | |
| add.f32 %f84, %f83, %f70; | |
| add.f32 %f85, %f84, %f72; | |
| max.f32 %f86, %f85, %f55; | |
| add.f32 %f87, %f57, 0f00000000; | |
| add.f32 %f88, %f87, %f59; | |
| add.f32 %f89, %f88, %f61; | |
| add.f32 %f90, %f89, %f63; | |
| add.f32 %f91, %f90, %f65; | |
| add.f32 %f92, %f91, %f67; | |
| add.f32 %f93, %f92, %f69; | |
| add.f32 %f94, %f93, %f71; | |
| div.rn.f32 %f95, %f94, %f86; | |
| mul.f32 %f96, %f6, %f95; | |
| sub.f32 %f97, %f55, %f96; | |
| max.f32 %f98, %f53, %f97; | |
| add.f32 %f8, %f77, %f98; | |
| mov.b32 %f99, %r5; | |
| div.rn.f32 %f100, %f99, %f56; | |
| add.f32 %f101, %f100, %f15; | |
| div.rn.f32 %f102, %f5, %f101; | |
| mul.f32 %f9, %f102, 0f3D800000; | |
| abs.f32 %f10, %f9; | |
| setp.ltu.f32 %p27, %f10, 0f3F19999A; | |
| @%p27 bra $L__BB0_7; | |
| bra.uni $L__BB0_6; | |
| $L__BB0_7: | |
| mul.f32 %f111, %f9, %f9; | |
| mov.f32 %f112, 0fBD563CAE; | |
| mov.f32 %f113, 0f3C80F082; | |
| fma.rn.f32 %f114, %f113, %f111, %f112; | |
| mov.f32 %f115, 0f3E085941; | |
| fma.rn.f32 %f116, %f114, %f111, %f115; | |
| mov.f32 %f117, 0fBEAAA9ED; | |
| fma.rn.f32 %f118, %f116, %f111, %f117; | |
| fma.rn.f32 %f120, %f118, %f111, %f53; | |
| fma.rn.f32 %f134, %f120, %f9, %f9; | |
| bra.uni $L__BB0_8; | |
| $L__BB0_6: | |
| mul.f32 %f103, %f10, 0f4038AA3B; | |
| ex2.approx.ftz.f32 %f104, %f103; | |
| add.f32 %f105, %f104, 0f3F800000; | |
| rcp.approx.ftz.f32 %f107, %f105; | |
| mov.f32 %f108, 0fC0000000; | |
| fma.rn.f32 %f109, %f107, %f108, %f55; | |
| setp.ge.f32 %p28, %f10, 0f41102CB4; | |
| selp.f32 %f110, 0f3F800000, %f109, %p28; | |
| mov.b32 %r57, %f110; | |
| mov.b32 %r58, %f9; | |
| and.b32 %r59, %r58, -2147483648; | |
| or.b32 %r60, %r59, %r57; | |
| mov.b32 %f134, %r60; | |
| $L__BB0_8: | |
| mul.f32 %f121, %f134, 0fC1800000; | |
| mul.f32 %f122, %f121, %f14; | |
| mul.f32 %f123, %f8, 0fBF000000; | |
| fma.rn.f32 %f124, %f123, 0f3EFD70A4, 0f3F800000; | |
| mov.f32 %f125, 0f3F800000; | |
| mov.f32 %f126, 0f3C23D70A; | |
| max.f32 %f127, %f124, %f126; | |
| min.f32 %f128, %f127, %f125; | |
| fma.rn.f32 %f129, %f128, %f122, %f7; | |
| abs.f32 %f130, %f129; | |
| setp.lt.f32 %p29, %f130, 0f7F800000; | |
| selp.f32 %f131, %f129, %f7, %p29; | |
| st.global.f32 [%rd1], %f131; | |
| $L__BB0_9: | |
| ret; | |
| } | |
| // .globl qpola_kernel_fp16 | |
| .visible .entry qpola_kernel_fp16( | |
| .param .u64 qpola_kernel_fp16_param_0, | |
| .param .u64 qpola_kernel_fp16_param_1, | |
| .param .f32 qpola_kernel_fp16_param_2, | |
| .param .f32 qpola_kernel_fp16_param_3, | |
| .param .u32 qpola_kernel_fp16_param_4 | |
| ) | |
| { | |
| .reg .pred %p<30>; | |
| .reg .b16 %rs<4>; | |
| .reg .f32 %f<141>; | |
| .reg .b32 %r<61>; | |
| .reg .b64 %rd<9>; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction[32]; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active[32]; | |
| ld.param.u64 %rd3, [qpola_kernel_fp16_param_0]; | |
| ld.param.u64 %rd2, [qpola_kernel_fp16_param_1]; | |
| ld.param.f32 %f14, [qpola_kernel_fp16_param_2]; | |
| ld.param.f32 %f15, [qpola_kernel_fp16_param_3]; | |
| ld.param.u32 %r6, [qpola_kernel_fp16_param_4]; | |
| cvta.to.global.u64 %rd4, %rd3; | |
| mov.u32 %r7, %ntid.x; | |
| mov.u32 %r8, %ctaid.x; | |
| mov.u32 %r1, %tid.x; | |
| mad.lo.s32 %r2, %r8, %r7, %r1; | |
| setp.lt.s32 %p1, %r2, %r6; | |
| setp.ge.s32 %p3, %r2, %r6; | |
| mul.wide.s32 %rd5, %r2, 2; | |
| add.s64 %rd1, %rd4, %rd5; | |
| mov.f32 %f138, 0f00000000; | |
| mov.f32 %f139, %f138; | |
| @%p3 bra $L__BB1_2; | |
| cvta.to.global.u64 %rd6, %rd2; | |
| add.s64 %rd8, %rd6, %rd5; | |
| ld.global.u16 %rs1, [%rd8]; | |
| // begin inline asm | |
| { cvt.f32.f16 %f138, %rs1;} | |
| // end inline asm | |
| ld.global.u16 %rs2, [%rd1]; | |
| // begin inline asm | |
| { cvt.f32.f16 %f139, %rs2;} | |
| // end inline asm | |
| $L__BB1_2: | |
| selp.f32 %f20, 0f3F800000, 0f00000000, %p1; | |
| abs.f32 %f21, %f138; | |
| setp.lt.f32 %p4, %f21, 0f7F800000; | |
| selp.f32 %f5, %f138, 0f00000000, %p4; | |
| setp.gt.f32 %p5, %f5, 0f00000000; | |
| setp.lt.f32 %p6, %f5, 0f00000000; | |
| selp.f32 %f22, 0fBF800000, 0f00000000, %p6; | |
| selp.f32 %f6, 0f3F800000, %f22, %p5; | |
| abs.f32 %f23, %f5; | |
| mov.b32 %r9, %f20; | |
| mov.u32 %r10, 2; | |
| mov.u32 %r11, 31; | |
| mov.u32 %r12, 16; | |
| mov.u32 %r13, -1; | |
| shfl.sync.down.b32 %r14|%p7, %r9, %r12, %r11, %r13; | |
| mov.b32 %f24, %r14; | |
| add.f32 %f25, %f20, %f24; | |
| mov.b32 %r15, %f6; | |
| shfl.sync.down.b32 %r16|%p8, %r15, %r12, %r11, %r13; | |
| mov.b32 %f26, %r16; | |
| add.f32 %f27, %f6, %f26; | |
| mov.b32 %r17, %f23; | |
| shfl.sync.down.b32 %r18|%p9, %r17, %r12, %r11, %r13; | |
| mov.b32 %f28, %r18; | |
| add.f32 %f29, %f23, %f28; | |
| mov.b32 %r19, %f25; | |
| mov.u32 %r20, 8; | |
| shfl.sync.down.b32 %r21|%p10, %r19, %r20, %r11, %r13; | |
| mov.b32 %f30, %r21; | |
| add.f32 %f31, %f25, %f30; | |
| mov.b32 %r22, %f27; | |
| shfl.sync.down.b32 %r23|%p11, %r22, %r20, %r11, %r13; | |
| mov.b32 %f32, %r23; | |
| add.f32 %f33, %f27, %f32; | |
| mov.b32 %r24, %f29; | |
| shfl.sync.down.b32 %r25|%p12, %r24, %r20, %r11, %r13; | |
| mov.b32 %f34, %r25; | |
| add.f32 %f35, %f29, %f34; | |
| mov.b32 %r26, %f31; | |
| mov.u32 %r27, 4; | |
| shfl.sync.down.b32 %r28|%p13, %r26, %r27, %r11, %r13; | |
| mov.b32 %f36, %r28; | |
| add.f32 %f37, %f31, %f36; | |
| mov.b32 %r29, %f33; | |
| shfl.sync.down.b32 %r30|%p14, %r29, %r27, %r11, %r13; | |
| mov.b32 %f38, %r30; | |
| add.f32 %f39, %f33, %f38; | |
| mov.b32 %r31, %f35; | |
| shfl.sync.down.b32 %r32|%p15, %r31, %r27, %r11, %r13; | |
| mov.b32 %f40, %r32; | |
| add.f32 %f41, %f35, %f40; | |
| mov.b32 %r33, %f37; | |
| shfl.sync.down.b32 %r34|%p16, %r33, %r10, %r11, %r13; | |
| mov.b32 %f42, %r34; | |
| add.f32 %f43, %f37, %f42; | |
| mov.b32 %r35, %f39; | |
| shfl.sync.down.b32 %r36|%p17, %r35, %r10, %r11, %r13; | |
| mov.b32 %f44, %r36; | |
| add.f32 %f45, %f39, %f44; | |
| mov.b32 %r37, %f41; | |
| shfl.sync.down.b32 %r38|%p18, %r37, %r10, %r11, %r13; | |
| mov.b32 %f46, %r38; | |
| add.f32 %f47, %f41, %f46; | |
| mov.b32 %r39, %f43; | |
| mov.u32 %r40, 1; | |
| shfl.sync.down.b32 %r41|%p19, %r39, %r40, %r11, %r13; | |
| mov.b32 %f48, %r41; | |
| add.f32 %f49, %f43, %f48; | |
| mov.b32 %r42, %f45; | |
| shfl.sync.down.b32 %r43|%p20, %r42, %r40, %r11, %r13; | |
| mov.b32 %f50, %r43; | |
| add.f32 %f51, %f45, %f50; | |
| mov.b32 %r44, %f47; | |
| shfl.sync.down.b32 %r45|%p21, %r44, %r40, %r11, %r13; | |
| mov.b32 %f52, %r45; | |
| add.f32 %f53, %f47, %f52; | |
| mov.b32 %r46, %f49; | |
| mov.u32 %r47, 0; | |
| shfl.sync.idx.b32 %r3|%p22, %r46, %r47, %r11, %r13; | |
| mov.b32 %r48, %f51; | |
| shfl.sync.idx.b32 %r4|%p23, %r48, %r47, %r11, %r13; | |
| mov.b32 %r49, %f53; | |
| shfl.sync.idx.b32 %r5|%p2, %r49, %r47, %r11, %r13; | |
| and.b32 %r50, %r1, 31; | |
| setp.ne.s32 %p24, %r50, 0; | |
| @%p24 bra $L__BB1_4; | |
| shr.u32 %r51, %r1, 3; | |
| and.b32 %r52, %r51, 536870908; | |
| mov.u32 %r53, _ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction; | |
| add.s32 %r54, %r53, %r52; | |
| st.shared.u32 [%r54], %r4; | |
| mov.u32 %r55, _ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active; | |
| add.s32 %r56, %r55, %r52; | |
| st.shared.u32 [%r56], %r3; | |
| $L__BB1_4: | |
| bar.sync 0; | |
| @%p3 bra $L__BB1_9; | |
| abs.f32 %f54, %f139; | |
| setp.lt.f32 %p26, %f54, 0f7F800000; | |
| selp.f32 %f7, %f139, 0f00000000, %p26; | |
| mov.f32 %f55, 0f00000000; | |
| mov.b32 %f56, %r3; | |
| mov.f32 %f57, 0f3F800000; | |
| max.f32 %f58, %f56, %f57; | |
| ld.shared.f32 %f59, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction]; | |
| ld.shared.f32 %f60, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active]; | |
| ld.shared.f32 %f61, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+4]; | |
| ld.shared.f32 %f62, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+4]; | |
| ld.shared.f32 %f63, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+8]; | |
| ld.shared.f32 %f64, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+8]; | |
| ld.shared.f32 %f65, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+12]; | |
| ld.shared.f32 %f66, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+12]; | |
| ld.shared.f32 %f67, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+16]; | |
| ld.shared.f32 %f68, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+16]; | |
| ld.shared.f32 %f69, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+20]; | |
| ld.shared.f32 %f70, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+20]; | |
| ld.shared.f32 %f71, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+24]; | |
| ld.shared.f32 %f72, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+24]; | |
| ld.shared.f32 %f73, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE17s_block_direction+28]; | |
| ld.shared.f32 %f74, [_ZZ17qpola_kernel_implI6__halfEvPT_PKS1_ffiE14s_block_active+28]; | |
| mov.b32 %f75, %r4; | |
| div.rn.f32 %f76, %f75, %f58; | |
| mul.f32 %f77, %f6, %f76; | |
| sub.f32 %f78, %f57, %f77; | |
| max.f32 %f79, %f55, %f78; | |
| add.f32 %f80, %f60, 0f00000000; | |
| add.f32 %f81, %f80, %f62; | |
| add.f32 %f82, %f81, %f64; | |
| add.f32 %f83, %f82, %f66; | |
| add.f32 %f84, %f83, %f68; | |
| add.f32 %f85, %f84, %f70; | |
| add.f32 %f86, %f85, %f72; | |
| add.f32 %f87, %f86, %f74; | |
| max.f32 %f88, %f87, %f57; | |
| add.f32 %f89, %f59, 0f00000000; | |
| add.f32 %f90, %f89, %f61; | |
| add.f32 %f91, %f90, %f63; | |
| add.f32 %f92, %f91, %f65; | |
| add.f32 %f93, %f92, %f67; | |
| add.f32 %f94, %f93, %f69; | |
| add.f32 %f95, %f94, %f71; | |
| add.f32 %f96, %f95, %f73; | |
| div.rn.f32 %f97, %f96, %f88; | |
| mul.f32 %f98, %f6, %f97; | |
| sub.f32 %f99, %f57, %f98; | |
| max.f32 %f100, %f55, %f99; | |
| add.f32 %f8, %f79, %f100; | |
| mov.b32 %f101, %r5; | |
| div.rn.f32 %f102, %f101, %f58; | |
| add.f32 %f103, %f102, %f15; | |
| div.rn.f32 %f104, %f5, %f103; | |
| mul.f32 %f9, %f104, 0f3E000000; | |
| abs.f32 %f10, %f9; | |
| setp.ltu.f32 %p27, %f10, 0f3F19999A; | |
| @%p27 bra $L__BB1_7; | |
| bra.uni $L__BB1_6; | |
| $L__BB1_7: | |
| mul.f32 %f113, %f9, %f9; | |
| mov.f32 %f114, 0fBD563CAE; | |
| mov.f32 %f115, 0f3C80F082; | |
| fma.rn.f32 %f116, %f115, %f113, %f114; | |
| mov.f32 %f117, 0f3E085941; | |
| fma.rn.f32 %f118, %f116, %f113, %f117; | |
| mov.f32 %f119, 0fBEAAA9ED; | |
| fma.rn.f32 %f120, %f118, %f113, %f119; | |
| fma.rn.f32 %f122, %f120, %f113, %f55; | |
| fma.rn.f32 %f140, %f122, %f9, %f9; | |
| bra.uni $L__BB1_8; | |
| $L__BB1_6: | |
| mul.f32 %f105, %f10, 0f4038AA3B; | |
| ex2.approx.ftz.f32 %f106, %f105; | |
| add.f32 %f107, %f106, 0f3F800000; | |
| rcp.approx.ftz.f32 %f109, %f107; | |
| mov.f32 %f110, 0fC0000000; | |
| fma.rn.f32 %f111, %f109, %f110, %f57; | |
| setp.ge.f32 %p28, %f10, 0f41102CB4; | |
| selp.f32 %f112, 0f3F800000, %f111, %p28; | |
| mov.b32 %r57, %f112; | |
| mov.b32 %r58, %f9; | |
| and.b32 %r59, %r58, -2147483648; | |
| or.b32 %r60, %r59, %r57; | |
| mov.b32 %f140, %r60; | |
| $L__BB1_8: | |
| mul.f32 %f124, %f140, 0fC1000000; | |
| mul.f32 %f125, %f124, %f14; | |
| mul.f32 %f126, %f8, 0fBF000000; | |
| fma.rn.f32 %f127, %f126, 0f3EFD70A4, 0f3F800000; | |
| mov.f32 %f128, 0f3F800000; | |
| mov.f32 %f129, 0f3C23D70A; | |
| max.f32 %f130, %f127, %f129; | |
| min.f32 %f131, %f130, %f128; | |
| fma.rn.f32 %f132, %f131, %f125, %f7; | |
| abs.f32 %f133, %f132; | |
| setp.lt.f32 %p29, %f133, 0f7F800000; | |
| selp.f32 %f134, %f132, %f7, %p29; | |
| mov.f32 %f135, 0f477FE000; | |
| min.f32 %f136, %f135, %f134; | |
| mov.f32 %f137, 0fC77FE000; | |
| max.f32 %f123, %f137, %f136; | |
| // begin inline asm | |
| { cvt.rn.f16.f32 %rs3, %f123;} | |
| // end inline asm | |
| st.global.u16 [%rd1], %rs3; | |
| $L__BB1_9: | |
| ret; | |
| } | |
| // .globl qpola_kernel_bf16 | |
| .visible .entry qpola_kernel_bf16( | |
| .param .u64 qpola_kernel_bf16_param_0, | |
| .param .u64 qpola_kernel_bf16_param_1, | |
| .param .f32 qpola_kernel_bf16_param_2, | |
| .param .f32 qpola_kernel_bf16_param_3, | |
| .param .u32 qpola_kernel_bf16_param_4 | |
| ) | |
| { | |
| .reg .pred %p<36>; | |
| .reg .b16 %rs<9>; | |
| .reg .f32 %f<141>; | |
| .reg .b32 %r<66>; | |
| .reg .b64 %rd<9>; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction[32]; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active[32]; | |
| ld.param.u64 %rd3, [qpola_kernel_bf16_param_0]; | |
| ld.param.u64 %rd2, [qpola_kernel_bf16_param_1]; | |
| ld.param.f32 %f14, [qpola_kernel_bf16_param_2]; | |
| ld.param.f32 %f15, [qpola_kernel_bf16_param_3]; | |
| ld.param.u32 %r7, [qpola_kernel_bf16_param_4]; | |
| cvta.to.global.u64 %rd4, %rd3; | |
| mov.u32 %r8, %ntid.x; | |
| mov.u32 %r9, %ctaid.x; | |
| mov.u32 %r1, %tid.x; | |
| mad.lo.s32 %r2, %r9, %r8, %r1; | |
| setp.lt.s32 %p1, %r2, %r7; | |
| setp.ge.s32 %p3, %r2, %r7; | |
| mul.wide.s32 %rd5, %r2, 2; | |
| add.s64 %rd1, %rd4, %rd5; | |
| mov.f32 %f138, 0f00000000; | |
| mov.f32 %f139, %f138; | |
| @%p3 bra $L__BB2_2; | |
| cvta.to.global.u64 %rd6, %rd2; | |
| add.s64 %rd8, %rd6, %rd5; | |
| ld.global.u16 %rs4, [%rd8]; | |
| // begin inline asm | |
| { mov.b32 %f138, {0,%rs4};} | |
| // end inline asm | |
| ld.global.u16 %rs5, [%rd1]; | |
| // begin inline asm | |
| { mov.b32 %f139, {0,%rs5};} | |
| // end inline asm | |
| $L__BB2_2: | |
| selp.f32 %f20, 0f3F800000, 0f00000000, %p1; | |
| abs.f32 %f21, %f138; | |
| setp.lt.f32 %p4, %f21, 0f7F800000; | |
| selp.f32 %f5, %f138, 0f00000000, %p4; | |
| setp.gt.f32 %p5, %f5, 0f00000000; | |
| setp.lt.f32 %p6, %f5, 0f00000000; | |
| selp.f32 %f22, 0fBF800000, 0f00000000, %p6; | |
| selp.f32 %f6, 0f3F800000, %f22, %p5; | |
| abs.f32 %f23, %f5; | |
| mov.b32 %r10, %f20; | |
| mov.u32 %r11, 2; | |
| mov.u32 %r12, 31; | |
| mov.u32 %r13, 16; | |
| mov.u32 %r14, -1; | |
| shfl.sync.down.b32 %r15|%p7, %r10, %r13, %r12, %r14; | |
| mov.b32 %f24, %r15; | |
| add.f32 %f25, %f20, %f24; | |
| mov.b32 %r16, %f6; | |
| shfl.sync.down.b32 %r17|%p8, %r16, %r13, %r12, %r14; | |
| mov.b32 %f26, %r17; | |
| add.f32 %f27, %f6, %f26; | |
| mov.b32 %r18, %f23; | |
| shfl.sync.down.b32 %r19|%p9, %r18, %r13, %r12, %r14; | |
| mov.b32 %f28, %r19; | |
| add.f32 %f29, %f23, %f28; | |
| mov.b32 %r20, %f25; | |
| mov.u32 %r21, 8; | |
| shfl.sync.down.b32 %r22|%p10, %r20, %r21, %r12, %r14; | |
| mov.b32 %f30, %r22; | |
| add.f32 %f31, %f25, %f30; | |
| mov.b32 %r23, %f27; | |
| shfl.sync.down.b32 %r24|%p11, %r23, %r21, %r12, %r14; | |
| mov.b32 %f32, %r24; | |
| add.f32 %f33, %f27, %f32; | |
| mov.b32 %r25, %f29; | |
| shfl.sync.down.b32 %r26|%p12, %r25, %r21, %r12, %r14; | |
| mov.b32 %f34, %r26; | |
| add.f32 %f35, %f29, %f34; | |
| mov.b32 %r27, %f31; | |
| mov.u32 %r28, 4; | |
| shfl.sync.down.b32 %r29|%p13, %r27, %r28, %r12, %r14; | |
| mov.b32 %f36, %r29; | |
| add.f32 %f37, %f31, %f36; | |
| mov.b32 %r30, %f33; | |
| shfl.sync.down.b32 %r31|%p14, %r30, %r28, %r12, %r14; | |
| mov.b32 %f38, %r31; | |
| add.f32 %f39, %f33, %f38; | |
| mov.b32 %r32, %f35; | |
| shfl.sync.down.b32 %r33|%p15, %r32, %r28, %r12, %r14; | |
| mov.b32 %f40, %r33; | |
| add.f32 %f41, %f35, %f40; | |
| mov.b32 %r34, %f37; | |
| shfl.sync.down.b32 %r35|%p16, %r34, %r11, %r12, %r14; | |
| mov.b32 %f42, %r35; | |
| add.f32 %f43, %f37, %f42; | |
| mov.b32 %r36, %f39; | |
| shfl.sync.down.b32 %r37|%p17, %r36, %r11, %r12, %r14; | |
| mov.b32 %f44, %r37; | |
| add.f32 %f45, %f39, %f44; | |
| mov.b32 %r38, %f41; | |
| shfl.sync.down.b32 %r39|%p18, %r38, %r11, %r12, %r14; | |
| mov.b32 %f46, %r39; | |
| add.f32 %f47, %f41, %f46; | |
| mov.b32 %r40, %f43; | |
| mov.u32 %r41, 1; | |
| shfl.sync.down.b32 %r42|%p19, %r40, %r41, %r12, %r14; | |
| mov.b32 %f48, %r42; | |
| add.f32 %f49, %f43, %f48; | |
| mov.b32 %r43, %f45; | |
| shfl.sync.down.b32 %r44|%p20, %r43, %r41, %r12, %r14; | |
| mov.b32 %f50, %r44; | |
| add.f32 %f51, %f45, %f50; | |
| mov.b32 %r45, %f47; | |
| shfl.sync.down.b32 %r46|%p21, %r45, %r41, %r12, %r14; | |
| mov.b32 %f52, %r46; | |
| add.f32 %f53, %f47, %f52; | |
| mov.b32 %r47, %f49; | |
| mov.u32 %r48, 0; | |
| shfl.sync.idx.b32 %r3|%p22, %r47, %r48, %r12, %r14; | |
| mov.b32 %r49, %f51; | |
| shfl.sync.idx.b32 %r4|%p23, %r49, %r48, %r12, %r14; | |
| mov.b32 %r50, %f53; | |
| shfl.sync.idx.b32 %r5|%p2, %r50, %r48, %r12, %r14; | |
| and.b32 %r51, %r1, 31; | |
| setp.ne.s32 %p24, %r51, 0; | |
| @%p24 bra $L__BB2_4; | |
| shr.u32 %r52, %r1, 3; | |
| and.b32 %r53, %r52, 536870908; | |
| mov.u32 %r54, _ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction; | |
| add.s32 %r55, %r54, %r53; | |
| st.shared.u32 [%r55], %r4; | |
| mov.u32 %r56, _ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active; | |
| add.s32 %r57, %r56, %r53; | |
| st.shared.u32 [%r57], %r3; | |
| $L__BB2_4: | |
| bar.sync 0; | |
| @%p3 bra $L__BB2_12; | |
| abs.f32 %f54, %f139; | |
| setp.lt.f32 %p26, %f54, 0f7F800000; | |
| selp.f32 %f7, %f139, 0f00000000, %p26; | |
| mov.f32 %f55, 0f00000000; | |
| mov.b32 %f56, %r3; | |
| mov.f32 %f57, 0f3F800000; | |
| max.f32 %f58, %f56, %f57; | |
| ld.shared.f32 %f59, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction]; | |
| ld.shared.f32 %f60, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active]; | |
| ld.shared.f32 %f61, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+4]; | |
| ld.shared.f32 %f62, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+4]; | |
| ld.shared.f32 %f63, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+8]; | |
| ld.shared.f32 %f64, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+8]; | |
| ld.shared.f32 %f65, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+12]; | |
| ld.shared.f32 %f66, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+12]; | |
| ld.shared.f32 %f67, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+16]; | |
| ld.shared.f32 %f68, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+16]; | |
| ld.shared.f32 %f69, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+20]; | |
| ld.shared.f32 %f70, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+20]; | |
| ld.shared.f32 %f71, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+24]; | |
| ld.shared.f32 %f72, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+24]; | |
| ld.shared.f32 %f73, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE17s_block_direction+28]; | |
| ld.shared.f32 %f74, [_ZZ17qpola_kernel_implI13__nv_bfloat16EvPT_PKS1_ffiE14s_block_active+28]; | |
| mov.b32 %f75, %r4; | |
| div.rn.f32 %f76, %f75, %f58; | |
| mul.f32 %f77, %f6, %f76; | |
| sub.f32 %f78, %f57, %f77; | |
| max.f32 %f79, %f55, %f78; | |
| add.f32 %f80, %f60, 0f00000000; | |
| add.f32 %f81, %f80, %f62; | |
| add.f32 %f82, %f81, %f64; | |
| add.f32 %f83, %f82, %f66; | |
| add.f32 %f84, %f83, %f68; | |
| add.f32 %f85, %f84, %f70; | |
| add.f32 %f86, %f85, %f72; | |
| add.f32 %f87, %f86, %f74; | |
| max.f32 %f88, %f87, %f57; | |
| add.f32 %f89, %f59, 0f00000000; | |
| add.f32 %f90, %f89, %f61; | |
| add.f32 %f91, %f90, %f63; | |
| add.f32 %f92, %f91, %f65; | |
| add.f32 %f93, %f92, %f67; | |
| add.f32 %f94, %f93, %f69; | |
| add.f32 %f95, %f94, %f71; | |
| add.f32 %f96, %f95, %f73; | |
| div.rn.f32 %f97, %f96, %f88; | |
| mul.f32 %f98, %f6, %f97; | |
| sub.f32 %f99, %f57, %f98; | |
| max.f32 %f100, %f55, %f99; | |
| add.f32 %f8, %f79, %f100; | |
| mov.b32 %f101, %r5; | |
| div.rn.f32 %f102, %f101, %f58; | |
| add.f32 %f103, %f102, %f15; | |
| div.rn.f32 %f104, %f5, %f103; | |
| mul.f32 %f9, %f104, 0f3E000000; | |
| abs.f32 %f10, %f9; | |
| setp.ltu.f32 %p27, %f10, 0f3F19999A; | |
| @%p27 bra $L__BB2_7; | |
| bra.uni $L__BB2_6; | |
| $L__BB2_7: | |
| mul.f32 %f113, %f9, %f9; | |
| mov.f32 %f114, 0fBD563CAE; | |
| mov.f32 %f115, 0f3C80F082; | |
| fma.rn.f32 %f116, %f115, %f113, %f114; | |
| mov.f32 %f117, 0f3E085941; | |
| fma.rn.f32 %f118, %f116, %f113, %f117; | |
| mov.f32 %f119, 0fBEAAA9ED; | |
| fma.rn.f32 %f120, %f118, %f113, %f119; | |
| fma.rn.f32 %f122, %f120, %f113, %f55; | |
| fma.rn.f32 %f140, %f122, %f9, %f9; | |
| bra.uni $L__BB2_8; | |
| $L__BB2_6: | |
| mul.f32 %f105, %f10, 0f4038AA3B; | |
| ex2.approx.ftz.f32 %f106, %f105; | |
| add.f32 %f107, %f106, 0f3F800000; | |
| rcp.approx.ftz.f32 %f109, %f107; | |
| mov.f32 %f110, 0fC0000000; | |
| fma.rn.f32 %f111, %f109, %f110, %f57; | |
| setp.ge.f32 %p28, %f10, 0f41102CB4; | |
| selp.f32 %f112, 0f3F800000, %f111, %p28; | |
| mov.b32 %r58, %f112; | |
| mov.b32 %r59, %f9; | |
| and.b32 %r60, %r59, -2147483648; | |
| or.b32 %r61, %r60, %r58; | |
| mov.b32 %f140, %r61; | |
| $L__BB2_8: | |
| mul.f32 %f123, %f140, 0fC1000000; | |
| mul.f32 %f124, %f123, %f14; | |
| mul.f32 %f125, %f8, 0fBF000000; | |
| fma.rn.f32 %f126, %f125, 0f3EFD70A4, 0f3F800000; | |
| mov.f32 %f127, 0f3F800000; | |
| mov.f32 %f128, 0f3C23D70A; | |
| max.f32 %f129, %f126, %f128; | |
| min.f32 %f130, %f129, %f127; | |
| fma.rn.f32 %f131, %f130, %f124, %f7; | |
| abs.f32 %f132, %f131; | |
| setp.lt.f32 %p29, %f132, 0f7F800000; | |
| selp.f32 %f133, %f131, %f7, %p29; | |
| mov.f32 %f134, 0f477FE000; | |
| min.f32 %f135, %f134, %f133; | |
| mov.f32 %f136, 0fC77FE000; | |
| max.f32 %f137, %f136, %f135; | |
| mov.b32 %r62, %f137; | |
| and.b32 %r63, %r62, 2147483647; | |
| setp.gt.u32 %p30, %r63, 2139095040; | |
| shl.b32 %r64, %r62, 16; | |
| shr.u32 %r65, %r62, 16; | |
| cvt.u16.u32 %rs6, %r65; | |
| selp.b32 %r6, 0, %r64, %p30; | |
| selp.b16 %rs8, 32767, %rs6, %p30; | |
| setp.gt.u32 %p31, %r6, -2147483648; | |
| @%p31 bra $L__BB2_10; | |
| setp.ne.s32 %p32, %r6, -2147483648; | |
| and.b16 %rs7, %rs8, 1; | |
| setp.eq.b16 %p33, %rs7, 1; | |
| not.pred %p34, %p33; | |
| or.pred %p35, %p32, %p34; | |
| @%p35 bra $L__BB2_11; | |
| $L__BB2_10: | |
| add.s16 %rs8, %rs8, 1; | |
| $L__BB2_11: | |
| st.global.u16 [%rd1], %rs8; | |
| $L__BB2_12: | |
| ret; | |
| } | |
| // .globl qpola_kernel_int8 | |
| .visible .entry qpola_kernel_int8( | |
| .param .u64 qpola_kernel_int8_param_0, | |
| .param .u64 qpola_kernel_int8_param_1, | |
| .param .f32 qpola_kernel_int8_param_2, | |
| .param .f32 qpola_kernel_int8_param_3, | |
| .param .u32 qpola_kernel_int8_param_4 | |
| ) | |
| { | |
| .reg .pred %p<30>; | |
| .reg .b16 %rs<3>; | |
| .reg .f32 %f<144>; | |
| .reg .b32 %r<64>; | |
| .reg .b64 %rd<9>; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction[32]; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active[32]; | |
| ld.param.u64 %rd3, [qpola_kernel_int8_param_0]; | |
| ld.param.u64 %rd2, [qpola_kernel_int8_param_1]; | |
| ld.param.f32 %f14, [qpola_kernel_int8_param_2]; | |
| ld.param.f32 %f15, [qpola_kernel_int8_param_3]; | |
| ld.param.u32 %r6, [qpola_kernel_int8_param_4]; | |
| cvta.to.global.u64 %rd4, %rd3; | |
| mov.u32 %r7, %ntid.x; | |
| mov.u32 %r8, %ctaid.x; | |
| mov.u32 %r1, %tid.x; | |
| mad.lo.s32 %r2, %r8, %r7, %r1; | |
| setp.lt.s32 %p1, %r2, %r6; | |
| setp.ge.s32 %p3, %r2, %r6; | |
| cvt.s64.s32 %rd5, %r2; | |
| add.s64 %rd1, %rd4, %rd5; | |
| mov.f32 %f141, 0f00000000; | |
| mov.f32 %f142, %f141; | |
| @%p3 bra $L__BB3_2; | |
| cvta.to.global.u64 %rd6, %rd2; | |
| add.s64 %rd8, %rd6, %rd5; | |
| ld.global.s8 %rs1, [%rd8]; | |
| cvt.rn.f32.s16 %f141, %rs1; | |
| ld.global.s8 %rs2, [%rd1]; | |
| cvt.rn.f32.s16 %f142, %rs2; | |
| $L__BB3_2: | |
| selp.f32 %f18, 0f3F800000, 0f00000000, %p1; | |
| abs.f32 %f19, %f141; | |
| setp.lt.f32 %p4, %f19, 0f7F800000; | |
| selp.f32 %f5, %f141, 0f00000000, %p4; | |
| setp.gt.f32 %p5, %f5, 0f00000000; | |
| setp.lt.f32 %p6, %f5, 0f00000000; | |
| selp.f32 %f20, 0fBF800000, 0f00000000, %p6; | |
| selp.f32 %f6, 0f3F800000, %f20, %p5; | |
| abs.f32 %f21, %f5; | |
| mov.b32 %r9, %f18; | |
| mov.u32 %r10, 2; | |
| mov.u32 %r11, 31; | |
| mov.u32 %r12, 16; | |
| mov.u32 %r13, -1; | |
| shfl.sync.down.b32 %r14|%p7, %r9, %r12, %r11, %r13; | |
| mov.b32 %f22, %r14; | |
| add.f32 %f23, %f18, %f22; | |
| mov.b32 %r15, %f6; | |
| shfl.sync.down.b32 %r16|%p8, %r15, %r12, %r11, %r13; | |
| mov.b32 %f24, %r16; | |
| add.f32 %f25, %f6, %f24; | |
| mov.b32 %r17, %f21; | |
| shfl.sync.down.b32 %r18|%p9, %r17, %r12, %r11, %r13; | |
| mov.b32 %f26, %r18; | |
| add.f32 %f27, %f21, %f26; | |
| mov.b32 %r19, %f23; | |
| mov.u32 %r20, 8; | |
| shfl.sync.down.b32 %r21|%p10, %r19, %r20, %r11, %r13; | |
| mov.b32 %f28, %r21; | |
| add.f32 %f29, %f23, %f28; | |
| mov.b32 %r22, %f25; | |
| shfl.sync.down.b32 %r23|%p11, %r22, %r20, %r11, %r13; | |
| mov.b32 %f30, %r23; | |
| add.f32 %f31, %f25, %f30; | |
| mov.b32 %r24, %f27; | |
| shfl.sync.down.b32 %r25|%p12, %r24, %r20, %r11, %r13; | |
| mov.b32 %f32, %r25; | |
| add.f32 %f33, %f27, %f32; | |
| mov.b32 %r26, %f29; | |
| mov.u32 %r27, 4; | |
| shfl.sync.down.b32 %r28|%p13, %r26, %r27, %r11, %r13; | |
| mov.b32 %f34, %r28; | |
| add.f32 %f35, %f29, %f34; | |
| mov.b32 %r29, %f31; | |
| shfl.sync.down.b32 %r30|%p14, %r29, %r27, %r11, %r13; | |
| mov.b32 %f36, %r30; | |
| add.f32 %f37, %f31, %f36; | |
| mov.b32 %r31, %f33; | |
| shfl.sync.down.b32 %r32|%p15, %r31, %r27, %r11, %r13; | |
| mov.b32 %f38, %r32; | |
| add.f32 %f39, %f33, %f38; | |
| mov.b32 %r33, %f35; | |
| shfl.sync.down.b32 %r34|%p16, %r33, %r10, %r11, %r13; | |
| mov.b32 %f40, %r34; | |
| add.f32 %f41, %f35, %f40; | |
| mov.b32 %r35, %f37; | |
| shfl.sync.down.b32 %r36|%p17, %r35, %r10, %r11, %r13; | |
| mov.b32 %f42, %r36; | |
| add.f32 %f43, %f37, %f42; | |
| mov.b32 %r37, %f39; | |
| shfl.sync.down.b32 %r38|%p18, %r37, %r10, %r11, %r13; | |
| mov.b32 %f44, %r38; | |
| add.f32 %f45, %f39, %f44; | |
| mov.b32 %r39, %f41; | |
| mov.u32 %r40, 1; | |
| shfl.sync.down.b32 %r41|%p19, %r39, %r40, %r11, %r13; | |
| mov.b32 %f46, %r41; | |
| add.f32 %f47, %f41, %f46; | |
| mov.b32 %r42, %f43; | |
| shfl.sync.down.b32 %r43|%p20, %r42, %r40, %r11, %r13; | |
| mov.b32 %f48, %r43; | |
| add.f32 %f49, %f43, %f48; | |
| mov.b32 %r44, %f45; | |
| shfl.sync.down.b32 %r45|%p21, %r44, %r40, %r11, %r13; | |
| mov.b32 %f50, %r45; | |
| add.f32 %f51, %f45, %f50; | |
| mov.b32 %r46, %f47; | |
| mov.u32 %r47, 0; | |
| shfl.sync.idx.b32 %r3|%p22, %r46, %r47, %r11, %r13; | |
| mov.b32 %r48, %f49; | |
| shfl.sync.idx.b32 %r4|%p23, %r48, %r47, %r11, %r13; | |
| mov.b32 %r49, %f51; | |
| shfl.sync.idx.b32 %r5|%p2, %r49, %r47, %r11, %r13; | |
| and.b32 %r50, %r1, 31; | |
| setp.ne.s32 %p24, %r50, 0; | |
| @%p24 bra $L__BB3_4; | |
| shr.u32 %r51, %r1, 3; | |
| and.b32 %r52, %r51, 536870908; | |
| mov.u32 %r53, _ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction; | |
| add.s32 %r54, %r53, %r52; | |
| st.shared.u32 [%r54], %r4; | |
| mov.u32 %r55, _ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active; | |
| add.s32 %r56, %r55, %r52; | |
| st.shared.u32 [%r56], %r3; | |
| $L__BB3_4: | |
| bar.sync 0; | |
| @%p3 bra $L__BB3_9; | |
| abs.f32 %f52, %f142; | |
| setp.lt.f32 %p26, %f52, 0f7F800000; | |
| selp.f32 %f7, %f142, 0f00000000, %p26; | |
| mov.f32 %f53, 0f00000000; | |
| mov.b32 %f54, %r3; | |
| mov.f32 %f55, 0f3F800000; | |
| max.f32 %f56, %f54, %f55; | |
| ld.shared.f32 %f57, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction]; | |
| ld.shared.f32 %f58, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active]; | |
| ld.shared.f32 %f59, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+4]; | |
| ld.shared.f32 %f60, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+4]; | |
| ld.shared.f32 %f61, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+8]; | |
| ld.shared.f32 %f62, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+8]; | |
| ld.shared.f32 %f63, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+12]; | |
| ld.shared.f32 %f64, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+12]; | |
| ld.shared.f32 %f65, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+16]; | |
| ld.shared.f32 %f66, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+16]; | |
| ld.shared.f32 %f67, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+20]; | |
| ld.shared.f32 %f68, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+20]; | |
| ld.shared.f32 %f69, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+24]; | |
| ld.shared.f32 %f70, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+24]; | |
| ld.shared.f32 %f71, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE17s_block_direction+28]; | |
| ld.shared.f32 %f72, [_ZZ17qpola_kernel_implIaEvPT_PKS0_ffiE14s_block_active+28]; | |
| mov.b32 %f73, %r4; | |
| div.rn.f32 %f74, %f73, %f56; | |
| mul.f32 %f75, %f6, %f74; | |
| sub.f32 %f76, %f55, %f75; | |
| max.f32 %f77, %f53, %f76; | |
| add.f32 %f78, %f58, 0f00000000; | |
| add.f32 %f79, %f78, %f60; | |
| add.f32 %f80, %f79, %f62; | |
| add.f32 %f81, %f80, %f64; | |
| add.f32 %f82, %f81, %f66; | |
| add.f32 %f83, %f82, %f68; | |
| add.f32 %f84, %f83, %f70; | |
| add.f32 %f85, %f84, %f72; | |
| max.f32 %f86, %f85, %f55; | |
| add.f32 %f87, %f57, 0f00000000; | |
| add.f32 %f88, %f87, %f59; | |
| add.f32 %f89, %f88, %f61; | |
| add.f32 %f90, %f89, %f63; | |
| add.f32 %f91, %f90, %f65; | |
| add.f32 %f92, %f91, %f67; | |
| add.f32 %f93, %f92, %f69; | |
| add.f32 %f94, %f93, %f71; | |
| div.rn.f32 %f95, %f94, %f86; | |
| mul.f32 %f96, %f6, %f95; | |
| sub.f32 %f97, %f55, %f96; | |
| max.f32 %f98, %f53, %f97; | |
| add.f32 %f8, %f77, %f98; | |
| mov.b32 %f99, %r5; | |
| div.rn.f32 %f100, %f99, %f56; | |
| add.f32 %f101, %f100, %f15; | |
| div.rn.f32 %f102, %f5, %f101; | |
| mul.f32 %f9, %f102, 0f3F000000; | |
| abs.f32 %f10, %f9; | |
| setp.ltu.f32 %p27, %f10, 0f3F19999A; | |
| @%p27 bra $L__BB3_7; | |
| bra.uni $L__BB3_6; | |
| $L__BB3_7: | |
| mul.f32 %f111, %f9, %f9; | |
| mov.f32 %f112, 0fBD563CAE; | |
| mov.f32 %f113, 0f3C80F082; | |
| fma.rn.f32 %f114, %f113, %f111, %f112; | |
| mov.f32 %f115, 0f3E085941; | |
| fma.rn.f32 %f116, %f114, %f111, %f115; | |
| mov.f32 %f117, 0fBEAAA9ED; | |
| fma.rn.f32 %f118, %f116, %f111, %f117; | |
| fma.rn.f32 %f120, %f118, %f111, %f53; | |
| fma.rn.f32 %f143, %f120, %f9, %f9; | |
| bra.uni $L__BB3_8; | |
| $L__BB3_6: | |
| mul.f32 %f103, %f10, 0f4038AA3B; | |
| ex2.approx.ftz.f32 %f104, %f103; | |
| add.f32 %f105, %f104, 0f3F800000; | |
| rcp.approx.ftz.f32 %f107, %f105; | |
| mov.f32 %f108, 0fC0000000; | |
| fma.rn.f32 %f109, %f107, %f108, %f55; | |
| setp.ge.f32 %p28, %f10, 0f41102CB4; | |
| selp.f32 %f110, 0f3F800000, %f109, %p28; | |
| mov.b32 %r57, %f110; | |
| mov.b32 %r58, %f9; | |
| and.b32 %r59, %r58, -2147483648; | |
| or.b32 %r60, %r59, %r57; | |
| mov.b32 %f143, %r60; | |
| $L__BB3_8: | |
| mul.f32 %f121, %f8, 0f3F000000; | |
| mul.f32 %f122, %f143, 0fC0000000; | |
| mul.f32 %f123, %f122, %f14; | |
| fma.rn.f32 %f124, %f121, 0fBEFD70A4, 0f3F800000; | |
| mov.f32 %f125, 0f3F800000; | |
| mov.f32 %f126, 0f3C23D70A; | |
| max.f32 %f127, %f124, %f126; | |
| min.f32 %f128, %f127, %f125; | |
| fma.rn.f32 %f129, %f128, %f123, %f7; | |
| abs.f32 %f130, %f129; | |
| setp.lt.f32 %p29, %f130, 0f7F800000; | |
| selp.f32 %f131, %f129, %f7, %p29; | |
| mov.f32 %f132, 0f42FE0000; | |
| min.f32 %f133, %f132, %f131; | |
| mov.f32 %f134, 0fC2FE0000; | |
| max.f32 %f135, %f134, %f133; | |
| fma.rn.f32 %f136, %f121, 0f3E4CCCCD, 0f3F800000; | |
| mul.f32 %f137, %f136, 0f3E800000; | |
| mad.lo.s32 %r61, %r2, 1664525, 7271263; | |
| and.b32 %r62, %r61, 16777215; | |
| cvt.rn.f32.u32 %f138, %r62; | |
| fma.rn.f32 %f139, %f138, 0f33800000, 0fBF000000; | |
| fma.rn.f32 %f140, %f139, %f137, %f135; | |
| cvt.rzi.s32.f32 %r63, %f140; | |
| st.global.u8 [%rd1], %r63; | |
| $L__BB3_9: | |
| ret; | |
| } | |
| // .globl qpola_kernel_fp8_e4m3 | |
| .visible .entry qpola_kernel_fp8_e4m3( | |
| .param .u64 qpola_kernel_fp8_e4m3_param_0, | |
| .param .u64 qpola_kernel_fp8_e4m3_param_1, | |
| .param .f32 qpola_kernel_fp8_e4m3_param_2, | |
| .param .f32 qpola_kernel_fp8_e4m3_param_3, | |
| .param .u32 qpola_kernel_fp8_e4m3_param_4 | |
| ) | |
| { | |
| .reg .pred %p<65>; | |
| .reg .b16 %rs<74>; | |
| .reg .f32 %f<144>; | |
| .reg .b32 %r<153>; | |
| .reg .f64 %fd<2>; | |
| .reg .b64 %rd<23>; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction[32]; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active[32]; | |
| ld.param.u64 %rd7, [qpola_kernel_fp8_e4m3_param_0]; | |
| ld.param.u64 %rd6, [qpola_kernel_fp8_e4m3_param_1]; | |
| ld.param.f32 %f14, [qpola_kernel_fp8_e4m3_param_2]; | |
| ld.param.f32 %f15, [qpola_kernel_fp8_e4m3_param_3]; | |
| ld.param.u32 %r44, [qpola_kernel_fp8_e4m3_param_4]; | |
| cvta.to.global.u64 %rd8, %rd7; | |
| mov.u32 %r45, %ntid.x; | |
| mov.u32 %r46, %ctaid.x; | |
| mov.u32 %r1, %tid.x; | |
| mad.lo.s32 %r2, %r46, %r45, %r1; | |
| setp.lt.s32 %p1, %r2, %r44; | |
| setp.ge.s32 %p3, %r2, %r44; | |
| cvt.s64.s32 %rd9, %r2; | |
| add.s64 %rd1, %rd8, %rd9; | |
| mov.f32 %f141, 0f00000000; | |
| mov.f32 %f142, %f141; | |
| @%p3 bra $L__BB4_36; | |
| cvta.to.global.u64 %rd10, %rd6; | |
| add.s64 %rd12, %rd10, %rd9; | |
| ld.global.u8 %rs1, [%rd12]; | |
| shl.b16 %rs36, %rs1, 8; | |
| and.b16 %rs2, %rs36, -32768; | |
| and.b16 %rs3, %rs36, 30720; | |
| shr.u16 %rs37, %rs3, 1; | |
| add.s16 %rs62, %rs37, 8192; | |
| and.b16 %rs63, %rs36, 1792; | |
| and.b16 %rs38, %rs1, 127; | |
| setp.eq.s16 %p4, %rs38, 127; | |
| mov.u16 %rs66, 32767; | |
| @%p4 bra $L__BB4_10; | |
| setp.eq.s16 %p5, %rs3, 0; | |
| @%p5 bra $L__BB4_4; | |
| shr.u16 %rs39, %rs63, 1; | |
| or.b16 %rs40, %rs39, %rs2; | |
| or.b16 %rs66, %rs40, %rs62; | |
| bra.uni $L__BB4_10; | |
| $L__BB4_4: | |
| setp.eq.s16 %p6, %rs63, 0; | |
| mov.u16 %rs64, 0; | |
| mov.u16 %rs65, %rs64; | |
| @%p6 bra $L__BB4_9; | |
| and.b16 %rs43, %rs1, 4; | |
| setp.ne.s16 %p7, %rs43, 0; | |
| @%p7 bra $L__BB4_8; | |
| mov.u16 %rs60, %rs63; | |
| $L__BB4_7: | |
| shl.b16 %rs63, %rs60, 1; | |
| add.s16 %rs62, %rs62, -1024; | |
| and.b16 %rs44, %rs60, 512; | |
| setp.eq.s16 %p8, %rs44, 0; | |
| mov.u16 %rs60, %rs63; | |
| @%p8 bra $L__BB4_7; | |
| $L__BB4_8: | |
| and.b16 %rs65, %rs63, 1022; | |
| mov.u16 %rs64, %rs62; | |
| $L__BB4_9: | |
| or.b16 %rs45, %rs64, %rs2; | |
| or.b16 %rs66, %rs45, %rs65; | |
| $L__BB4_10: | |
| cvt.u32.u16 %r47, %rs66; | |
| shr.u32 %r144, %r47, 15; | |
| bfe.u32 %r4, %r47, 10, 5; | |
| and.b32 %r5, %r47, 1023; | |
| mul.wide.u16 %r48, %rs66, 8192; | |
| and.b32 %r143, %r48, 8380416; | |
| setp.eq.s32 %p9, %r4, 31; | |
| @%p9 bra $L__BB4_17; | |
| bra.uni $L__BB4_11; | |
| $L__BB4_17: | |
| setp.eq.s32 %p13, %r5, 0; | |
| selp.b32 %r144, %r144, 0, %p13; | |
| selp.b32 %r146, 0, 8388607, %p13; | |
| mov.u32 %r145, 255; | |
| bra.uni $L__BB4_18; | |
| $L__BB4_11: | |
| setp.eq.s32 %p10, %r4, 0; | |
| @%p10 bra $L__BB4_13; | |
| add.s32 %r145, %r4, 112; | |
| mov.u32 %r146, %r143; | |
| bra.uni $L__BB4_18; | |
| $L__BB4_13: | |
| setp.eq.s32 %p11, %r5, 0; | |
| mov.u32 %r145, 0; | |
| mov.u32 %r146, %r145; | |
| @%p11 bra $L__BB4_18; | |
| mov.u32 %r145, 113; | |
| $L__BB4_15: | |
| shl.b32 %r10, %r143, 1; | |
| add.s32 %r145, %r145, -1; | |
| and.b32 %r52, %r143, 4194304; | |
| setp.eq.s32 %p12, %r52, 0; | |
| mov.u32 %r143, %r10; | |
| @%p12 bra $L__BB4_15; | |
| and.b32 %r146, %r10, 8388604; | |
| $L__BB4_18: | |
| shl.b32 %r54, %r145, 23; | |
| shl.b32 %r55, %r144, 31; | |
| or.b32 %r56, %r54, %r55; | |
| or.b32 %r18, %r56, %r146; | |
| ld.global.u8 %rs18, [%rd1]; | |
| shl.b16 %rs47, %rs18, 8; | |
| and.b16 %rs19, %rs47, -32768; | |
| and.b16 %rs20, %rs47, 30720; | |
| shr.u16 %rs48, %rs20, 1; | |
| add.s16 %rs69, %rs48, 8192; | |
| and.b16 %rs70, %rs47, 1792; | |
| and.b16 %rs49, %rs18, 127; | |
| setp.eq.s16 %p14, %rs49, 127; | |
| mov.u16 %rs73, 32767; | |
| @%p14 bra $L__BB4_27; | |
| setp.eq.s16 %p15, %rs20, 0; | |
| @%p15 bra $L__BB4_21; | |
| shr.u16 %rs50, %rs70, 1; | |
| or.b16 %rs51, %rs50, %rs19; | |
| or.b16 %rs73, %rs51, %rs69; | |
| bra.uni $L__BB4_27; | |
| $L__BB4_21: | |
| setp.eq.s16 %p16, %rs70, 0; | |
| mov.u16 %rs71, 0; | |
| mov.u16 %rs72, %rs71; | |
| @%p16 bra $L__BB4_26; | |
| and.b16 %rs54, %rs18, 4; | |
| setp.ne.s16 %p17, %rs54, 0; | |
| @%p17 bra $L__BB4_25; | |
| mov.u16 %rs67, %rs70; | |
| $L__BB4_24: | |
| shl.b16 %rs70, %rs67, 1; | |
| add.s16 %rs69, %rs69, -1024; | |
| and.b16 %rs55, %rs67, 512; | |
| setp.eq.s16 %p18, %rs55, 0; | |
| mov.u16 %rs67, %rs70; | |
| @%p18 bra $L__BB4_24; | |
| $L__BB4_25: | |
| and.b16 %rs72, %rs70, 1022; | |
| mov.u16 %rs71, %rs69; | |
| $L__BB4_26: | |
| or.b16 %rs56, %rs71, %rs19; | |
| or.b16 %rs73, %rs56, %rs72; | |
| $L__BB4_27: | |
| cvt.u32.u16 %r57, %rs73; | |
| shr.u32 %r149, %r57, 15; | |
| bfe.u32 %r20, %r57, 10, 5; | |
| and.b32 %r21, %r57, 1023; | |
| mul.wide.u16 %r58, %rs73, 8192; | |
| and.b32 %r148, %r58, 8380416; | |
| setp.eq.s32 %p19, %r20, 31; | |
| @%p19 bra $L__BB4_34; | |
| bra.uni $L__BB4_28; | |
| $L__BB4_34: | |
| setp.eq.s32 %p23, %r21, 0; | |
| selp.b32 %r149, %r149, 0, %p23; | |
| selp.b32 %r151, 0, 8388607, %p23; | |
| mov.u32 %r150, 255; | |
| bra.uni $L__BB4_35; | |
| $L__BB4_28: | |
| setp.eq.s32 %p20, %r20, 0; | |
| @%p20 bra $L__BB4_30; | |
| add.s32 %r150, %r20, 112; | |
| mov.u32 %r151, %r148; | |
| bra.uni $L__BB4_35; | |
| $L__BB4_30: | |
| setp.eq.s32 %p21, %r21, 0; | |
| mov.u32 %r150, 0; | |
| mov.u32 %r151, %r150; | |
| @%p21 bra $L__BB4_35; | |
| mov.u32 %r150, 113; | |
| $L__BB4_32: | |
| shl.b32 %r26, %r148, 1; | |
| add.s32 %r150, %r150, -1; | |
| and.b32 %r62, %r148, 4194304; | |
| setp.eq.s32 %p22, %r62, 0; | |
| mov.u32 %r148, %r26; | |
| @%p22 bra $L__BB4_32; | |
| and.b32 %r151, %r26, 8388604; | |
| $L__BB4_35: | |
| shl.b32 %r64, %r150, 23; | |
| shl.b32 %r65, %r149, 31; | |
| or.b32 %r66, %r64, %r65; | |
| or.b32 %r67, %r66, %r151; | |
| mov.b32 %f142, %r67; | |
| mov.b32 %f141, %r18; | |
| $L__BB4_36: | |
| selp.f32 %f18, 0f3F800000, 0f00000000, %p1; | |
| abs.f32 %f19, %f141; | |
| setp.lt.f32 %p24, %f19, 0f7F800000; | |
| selp.f32 %f5, %f141, 0f00000000, %p24; | |
| setp.gt.f32 %p25, %f5, 0f00000000; | |
| setp.lt.f32 %p26, %f5, 0f00000000; | |
| selp.f32 %f20, 0fBF800000, 0f00000000, %p26; | |
| selp.f32 %f6, 0f3F800000, %f20, %p25; | |
| abs.f32 %f21, %f5; | |
| mov.b32 %r68, %f18; | |
| mov.u32 %r69, 2; | |
| mov.u32 %r70, 31; | |
| mov.u32 %r71, 16; | |
| mov.u32 %r72, -1; | |
| shfl.sync.down.b32 %r73|%p27, %r68, %r71, %r70, %r72; | |
| mov.b32 %f22, %r73; | |
| add.f32 %f23, %f18, %f22; | |
| mov.b32 %r74, %f6; | |
| shfl.sync.down.b32 %r75|%p28, %r74, %r71, %r70, %r72; | |
| mov.b32 %f24, %r75; | |
| add.f32 %f25, %f6, %f24; | |
| mov.b32 %r76, %f21; | |
| shfl.sync.down.b32 %r77|%p29, %r76, %r71, %r70, %r72; | |
| mov.b32 %f26, %r77; | |
| add.f32 %f27, %f21, %f26; | |
| mov.b32 %r78, %f23; | |
| mov.u32 %r79, 8; | |
| shfl.sync.down.b32 %r80|%p30, %r78, %r79, %r70, %r72; | |
| mov.b32 %f28, %r80; | |
| add.f32 %f29, %f23, %f28; | |
| mov.b32 %r81, %f25; | |
| shfl.sync.down.b32 %r82|%p31, %r81, %r79, %r70, %r72; | |
| mov.b32 %f30, %r82; | |
| add.f32 %f31, %f25, %f30; | |
| mov.b32 %r83, %f27; | |
| shfl.sync.down.b32 %r84|%p32, %r83, %r79, %r70, %r72; | |
| mov.b32 %f32, %r84; | |
| add.f32 %f33, %f27, %f32; | |
| mov.b32 %r85, %f29; | |
| mov.u32 %r86, 4; | |
| shfl.sync.down.b32 %r87|%p33, %r85, %r86, %r70, %r72; | |
| mov.b32 %f34, %r87; | |
| add.f32 %f35, %f29, %f34; | |
| mov.b32 %r88, %f31; | |
| shfl.sync.down.b32 %r89|%p34, %r88, %r86, %r70, %r72; | |
| mov.b32 %f36, %r89; | |
| add.f32 %f37, %f31, %f36; | |
| mov.b32 %r90, %f33; | |
| shfl.sync.down.b32 %r91|%p35, %r90, %r86, %r70, %r72; | |
| mov.b32 %f38, %r91; | |
| add.f32 %f39, %f33, %f38; | |
| mov.b32 %r92, %f35; | |
| shfl.sync.down.b32 %r93|%p36, %r92, %r69, %r70, %r72; | |
| mov.b32 %f40, %r93; | |
| add.f32 %f41, %f35, %f40; | |
| mov.b32 %r94, %f37; | |
| shfl.sync.down.b32 %r95|%p37, %r94, %r69, %r70, %r72; | |
| mov.b32 %f42, %r95; | |
| add.f32 %f43, %f37, %f42; | |
| mov.b32 %r96, %f39; | |
| shfl.sync.down.b32 %r97|%p38, %r96, %r69, %r70, %r72; | |
| mov.b32 %f44, %r97; | |
| add.f32 %f45, %f39, %f44; | |
| mov.b32 %r98, %f41; | |
| mov.u32 %r99, 1; | |
| shfl.sync.down.b32 %r100|%p39, %r98, %r99, %r70, %r72; | |
| mov.b32 %f46, %r100; | |
| add.f32 %f47, %f41, %f46; | |
| mov.b32 %r101, %f43; | |
| shfl.sync.down.b32 %r102|%p40, %r101, %r99, %r70, %r72; | |
| mov.b32 %f48, %r102; | |
| add.f32 %f49, %f43, %f48; | |
| mov.b32 %r103, %f45; | |
| shfl.sync.down.b32 %r104|%p41, %r103, %r99, %r70, %r72; | |
| mov.b32 %f50, %r104; | |
| add.f32 %f51, %f45, %f50; | |
| mov.b32 %r105, %f47; | |
| mov.u32 %r106, 0; | |
| shfl.sync.idx.b32 %r34|%p42, %r105, %r106, %r70, %r72; | |
| mov.b32 %r107, %f49; | |
| shfl.sync.idx.b32 %r35|%p43, %r107, %r106, %r70, %r72; | |
| mov.b32 %r108, %f51; | |
| shfl.sync.idx.b32 %r36|%p2, %r108, %r106, %r70, %r72; | |
| and.b32 %r109, %r1, 31; | |
| setp.ne.s32 %p44, %r109, 0; | |
| @%p44 bra $L__BB4_38; | |
| shr.u32 %r110, %r1, 3; | |
| and.b32 %r111, %r110, 536870908; | |
| mov.u32 %r112, _ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction; | |
| add.s32 %r113, %r112, %r111; | |
| st.shared.u32 [%r113], %r35; | |
| mov.u32 %r114, _ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active; | |
| add.s32 %r115, %r114, %r111; | |
| st.shared.u32 [%r115], %r34; | |
| $L__BB4_38: | |
| bar.sync 0; | |
| @%p3 bra $L__BB4_51; | |
| abs.f32 %f52, %f142; | |
| setp.lt.f32 %p46, %f52, 0f7F800000; | |
| selp.f32 %f7, %f142, 0f00000000, %p46; | |
| mov.f32 %f53, 0f00000000; | |
| mov.b32 %f54, %r34; | |
| mov.f32 %f55, 0f3F800000; | |
| max.f32 %f56, %f54, %f55; | |
| ld.shared.f32 %f57, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction]; | |
| ld.shared.f32 %f58, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active]; | |
| ld.shared.f32 %f59, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+4]; | |
| ld.shared.f32 %f60, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+4]; | |
| ld.shared.f32 %f61, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+8]; | |
| ld.shared.f32 %f62, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+8]; | |
| ld.shared.f32 %f63, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+12]; | |
| ld.shared.f32 %f64, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+12]; | |
| ld.shared.f32 %f65, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+16]; | |
| ld.shared.f32 %f66, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+16]; | |
| ld.shared.f32 %f67, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+20]; | |
| ld.shared.f32 %f68, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+20]; | |
| ld.shared.f32 %f69, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+24]; | |
| ld.shared.f32 %f70, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+24]; | |
| ld.shared.f32 %f71, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE17s_block_direction+28]; | |
| ld.shared.f32 %f72, [_ZZ17qpola_kernel_implI13__nv_fp8_e4m3EvPT_PKS1_ffiE14s_block_active+28]; | |
| mov.b32 %f73, %r35; | |
| div.rn.f32 %f74, %f73, %f56; | |
| mul.f32 %f75, %f6, %f74; | |
| sub.f32 %f76, %f55, %f75; | |
| max.f32 %f77, %f53, %f76; | |
| add.f32 %f78, %f58, 0f00000000; | |
| add.f32 %f79, %f78, %f60; | |
| add.f32 %f80, %f79, %f62; | |
| add.f32 %f81, %f80, %f64; | |
| add.f32 %f82, %f81, %f66; | |
| add.f32 %f83, %f82, %f68; | |
| add.f32 %f84, %f83, %f70; | |
| add.f32 %f85, %f84, %f72; | |
| max.f32 %f86, %f85, %f55; | |
| add.f32 %f87, %f57, 0f00000000; | |
| add.f32 %f88, %f87, %f59; | |
| add.f32 %f89, %f88, %f61; | |
| add.f32 %f90, %f89, %f63; | |
| add.f32 %f91, %f90, %f65; | |
| add.f32 %f92, %f91, %f67; | |
| add.f32 %f93, %f92, %f69; | |
| add.f32 %f94, %f93, %f71; | |
| div.rn.f32 %f95, %f94, %f86; | |
| mul.f32 %f96, %f6, %f95; | |
| sub.f32 %f97, %f55, %f96; | |
| max.f32 %f98, %f53, %f97; | |
| add.f32 %f8, %f77, %f98; | |
| mov.b32 %f99, %r36; | |
| div.rn.f32 %f100, %f99, %f56; | |
| add.f32 %f101, %f100, %f15; | |
| div.rn.f32 %f102, %f5, %f101; | |
| mul.f32 %f9, %f102, 0f3E800000; | |
| abs.f32 %f10, %f9; | |
| setp.ltu.f32 %p47, %f10, 0f3F19999A; | |
| @%p47 bra $L__BB4_41; | |
| bra.uni $L__BB4_40; | |
| $L__BB4_41: | |
| mul.f32 %f111, %f9, %f9; | |
| mov.f32 %f112, 0fBD563CAE; | |
| mov.f32 %f113, 0f3C80F082; | |
| fma.rn.f32 %f114, %f113, %f111, %f112; | |
| mov.f32 %f115, 0f3E085941; | |
| fma.rn.f32 %f116, %f114, %f111, %f115; | |
| mov.f32 %f117, 0fBEAAA9ED; | |
| fma.rn.f32 %f118, %f116, %f111, %f117; | |
| fma.rn.f32 %f120, %f118, %f111, %f53; | |
| fma.rn.f32 %f143, %f120, %f9, %f9; | |
| bra.uni $L__BB4_42; | |
| $L__BB4_40: | |
| mul.f32 %f103, %f10, 0f4038AA3B; | |
| ex2.approx.ftz.f32 %f104, %f103; | |
| add.f32 %f105, %f104, 0f3F800000; | |
| rcp.approx.ftz.f32 %f107, %f105; | |
| mov.f32 %f108, 0fC0000000; | |
| fma.rn.f32 %f109, %f107, %f108, %f55; | |
| setp.ge.f32 %p48, %f10, 0f41102CB4; | |
| selp.f32 %f110, 0f3F800000, %f109, %p48; | |
| mov.b32 %r116, %f110; | |
| mov.b32 %r117, %f9; | |
| and.b32 %r118, %r117, -2147483648; | |
| or.b32 %r119, %r118, %r116; | |
| mov.b32 %f143, %r119; | |
| $L__BB4_42: | |
| mul.f32 %f121, %f8, 0f3F000000; | |
| mul.f32 %f122, %f143, 0fC0800000; | |
| mul.f32 %f123, %f122, %f14; | |
| fma.rn.f32 %f124, %f121, 0fBEFD70A4, 0f3F800000; | |
| mov.f32 %f125, 0f3F800000; | |
| mov.f32 %f126, 0f3C23D70A; | |
| max.f32 %f127, %f124, %f126; | |
| min.f32 %f128, %f127, %f125; | |
| fma.rn.f32 %f129, %f128, %f123, %f7; | |
| abs.f32 %f130, %f129; | |
| setp.lt.f32 %p49, %f130, 0f7F800000; | |
| selp.f32 %f131, %f129, %f7, %p49; | |
| mov.f32 %f132, 0f43E00000; | |
| min.f32 %f133, %f132, %f131; | |
| mov.f32 %f134, 0fC3E00000; | |
| max.f32 %f135, %f134, %f133; | |
| fma.rn.f32 %f136, %f121, 0f3E4CCCCD, 0f3F800000; | |
| mul.f32 %f137, %f136, 0f3C800000; | |
| mad.lo.s32 %r121, %r2, 1664525, 7271263; | |
| and.b32 %r122, %r121, 16777215; | |
| cvt.rn.f32.u32 %f138, %r122; | |
| fma.rn.f32 %f139, %f138, 0f33800000, 0fBF000000; | |
| fma.rn.f32 %f140, %f139, %f137, %f135; | |
| mov.b32 %r123, %f140; | |
| and.b32 %r124, %r123, 2147483647; | |
| setp.gt.u32 %p50, %r124, 2139095040; | |
| cvt.f64.f32 %fd1, %f140; | |
| mov.b64 %rd13, %fd1; | |
| selp.b64 %rd2, 9223372036317904896, %rd13, %p50; | |
| shr.u64 %rd14, %rd2, 52; | |
| cvt.u32.u64 %r125, %rd14; | |
| add.s32 %r37, %r125, 8; | |
| shr.u64 %rd15, %rd2, 49; | |
| cvt.u32.u64 %r38, %rd15; | |
| and.b32 %r39, %r38, 7; | |
| and.b64 %rd3, %rd2, 9223372036854775807; | |
| setp.lt.u64 %p51, %rd3, 4562146422526312449; | |
| mov.u32 %r152, 0; | |
| @%p51 bra $L__BB4_50; | |
| setp.gt.u64 %p52, %rd3, 9218868437227405312; | |
| mov.u32 %r152, 127; | |
| @%p52 bra $L__BB4_50; | |
| setp.gt.u64 %p53, %rd3, 4646870390516219904; | |
| mov.u32 %r152, 126; | |
| @%p53 bra $L__BB4_50; | |
| setp.lt.u64 %p54, %rd3, 4580160821035794432; | |
| @%p54 bra $L__BB4_47; | |
| bra.uni $L__BB4_46; | |
| $L__BB4_47: | |
| cvt.u16.u32 %rs57, %r37; | |
| mov.u16 %rs58, 1; | |
| sub.s16 %rs59, %rs58, %rs57; | |
| cvt.u32.u16 %r134, %rs59; | |
| and.b32 %r135, %r134, 255; | |
| or.b32 %r136, %r39, 8; | |
| shr.u32 %r152, %r136, %r135; | |
| mov.u64 %rd17, 562949953421312; | |
| shl.b64 %rd18, %rd17, %r135; | |
| add.s64 %rd19, %rd18, -1; | |
| or.b64 %rd20, %rd2, 4503599627370496; | |
| and.b64 %rd4, %rd19, %rd20; | |
| mov.u64 %rd21, 281474976710656; | |
| shl.b64 %rd5, %rd21, %r135; | |
| setp.gt.u64 %p60, %rd4, %rd5; | |
| @%p60 bra $L__BB4_49; | |
| setp.ne.s64 %p61, %rd4, %rd5; | |
| and.b32 %r137, %r152, 1; | |
| setp.eq.b32 %p62, %r137, 1; | |
| not.pred %p63, %p62; | |
| or.pred %p64, %p63, %p61; | |
| @%p64 bra $L__BB4_50; | |
| $L__BB4_49: | |
| and.b32 %r138, %r152, 255; | |
| add.s32 %r152, %r138, 1; | |
| bra.uni $L__BB4_50; | |
| $L__BB4_46: | |
| shl.b32 %r128, %r37, 3; | |
| and.b32 %r129, %r128, 2040; | |
| or.b32 %r130, %r129, %r39; | |
| and.b64 %rd16, %rd2, 562949953421311; | |
| setp.gt.u64 %p55, %rd16, 281474976710656; | |
| setp.eq.s64 %p56, %rd16, 281474976710656; | |
| and.b32 %r131, %r38, 1; | |
| setp.eq.b32 %p57, %r131, 1; | |
| and.pred %p58, %p56, %p57; | |
| or.pred %p59, %p55, %p58; | |
| and.b32 %r132, %r130, 255; | |
| add.s32 %r133, %r132, 1; | |
| selp.b32 %r152, %r133, %r130, %p59; | |
| $L__BB4_50: | |
| shr.u64 %rd22, %rd2, 63; | |
| cvt.u32.u64 %r139, %rd22; | |
| shl.b32 %r140, %r139, 7; | |
| or.b32 %r141, %r152, %r140; | |
| st.global.u8 [%rd1], %r141; | |
| $L__BB4_51: | |
| ret; | |
| } | |
| // .globl qpola_kernel_fp8_e5m2 | |
| .visible .entry qpola_kernel_fp8_e5m2( | |
| .param .u64 qpola_kernel_fp8_e5m2_param_0, | |
| .param .u64 qpola_kernel_fp8_e5m2_param_1, | |
| .param .f32 qpola_kernel_fp8_e5m2_param_2, | |
| .param .f32 qpola_kernel_fp8_e5m2_param_3, | |
| .param .u32 qpola_kernel_fp8_e5m2_param_4 | |
| ) | |
| { | |
| .reg .pred %p<57>; | |
| .reg .b16 %rs<12>; | |
| .reg .f32 %f<144>; | |
| .reg .b32 %r<153>; | |
| .reg .f64 %fd<2>; | |
| .reg .b64 %rd<23>; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction[32]; | |
| // demoted variable | |
| .shared .align 4 .b8 _ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active[32]; | |
| ld.param.u64 %rd7, [qpola_kernel_fp8_e5m2_param_0]; | |
| ld.param.u64 %rd6, [qpola_kernel_fp8_e5m2_param_1]; | |
| ld.param.f32 %f14, [qpola_kernel_fp8_e5m2_param_2]; | |
| ld.param.f32 %f15, [qpola_kernel_fp8_e5m2_param_3]; | |
| ld.param.u32 %r45, [qpola_kernel_fp8_e5m2_param_4]; | |
| cvta.to.global.u64 %rd8, %rd7; | |
| mov.u32 %r46, %ntid.x; | |
| mov.u32 %r47, %ctaid.x; | |
| mov.u32 %r1, %tid.x; | |
| mad.lo.s32 %r2, %r47, %r46, %r1; | |
| setp.lt.s32 %p1, %r2, %r45; | |
| setp.ge.s32 %p3, %r2, %r45; | |
| cvt.s64.s32 %rd9, %r2; | |
| add.s64 %rd1, %rd8, %rd9; | |
| mov.f32 %f141, 0f00000000; | |
| mov.f32 %f142, %f141; | |
| @%p3 bra $L__BB5_18; | |
| cvta.to.global.u64 %rd10, %rd6; | |
| add.s64 %rd12, %rd10, %rd9; | |
| ld.global.u8 %rs1, [%rd12]; | |
| shl.b16 %rs2, %rs1, 8; | |
| and.b16 %rs3, %rs1, 127; | |
| setp.gt.u16 %p4, %rs3, 124; | |
| selp.b16 %rs4, 32767, %rs2, %p4; | |
| cvt.u32.u16 %r48, %rs4; | |
| shr.u32 %r144, %r48, 15; | |
| bfe.u32 %r4, %r48, 10, 5; | |
| and.b32 %r5, %r48, 1023; | |
| mul.wide.u16 %r49, %rs4, 8192; | |
| and.b32 %r143, %r49, 8380416; | |
| setp.eq.s32 %p5, %r4, 31; | |
| @%p5 bra $L__BB5_8; | |
| bra.uni $L__BB5_2; | |
| $L__BB5_8: | |
| setp.eq.s32 %p9, %r5, 0; | |
| selp.b32 %r144, %r144, 0, %p9; | |
| selp.b32 %r146, 0, 8388607, %p9; | |
| mov.u32 %r145, 255; | |
| bra.uni $L__BB5_9; | |
| $L__BB5_2: | |
| setp.eq.s32 %p6, %r4, 0; | |
| @%p6 bra $L__BB5_4; | |
| add.s32 %r145, %r4, 112; | |
| mov.u32 %r146, %r143; | |
| bra.uni $L__BB5_9; | |
| $L__BB5_4: | |
| setp.eq.s32 %p7, %r5, 0; | |
| mov.u32 %r145, 0; | |
| mov.u32 %r146, %r145; | |
| @%p7 bra $L__BB5_9; | |
| mov.u32 %r145, 113; | |
| $L__BB5_6: | |
| shl.b32 %r10, %r143, 1; | |
| add.s32 %r145, %r145, -1; | |
| and.b32 %r53, %r143, 4194304; | |
| setp.eq.s32 %p8, %r53, 0; | |
| mov.u32 %r143, %r10; | |
| @%p8 bra $L__BB5_6; | |
| and.b32 %r146, %r10, 8388604; | |
| $L__BB5_9: | |
| shl.b32 %r55, %r145, 23; | |
| shl.b32 %r56, %r144, 31; | |
| or.b32 %r57, %r55, %r56; | |
| or.b32 %r18, %r57, %r146; | |
| ld.global.u8 %rs5, [%rd1]; | |
| shl.b16 %rs6, %rs5, 8; | |
| and.b16 %rs7, %rs5, 127; | |
| setp.gt.u16 %p10, %rs7, 124; | |
| selp.b16 %rs8, 32767, %rs6, %p10; | |
| cvt.u32.u16 %r58, %rs8; | |
| shr.u32 %r149, %r58, 15; | |
| bfe.u32 %r20, %r58, 10, 5; | |
| and.b32 %r21, %r58, 1023; | |
| mul.wide.u16 %r59, %rs8, 8192; | |
| and.b32 %r148, %r59, 8380416; | |
| setp.eq.s32 %p11, %r20, 31; | |
| @%p11 bra $L__BB5_16; | |
| bra.uni $L__BB5_10; | |
| $L__BB5_16: | |
| setp.eq.s32 %p15, %r21, 0; | |
| selp.b32 %r149, %r149, 0, %p15; | |
| selp.b32 %r151, 0, 8388607, %p15; | |
| mov.u32 %r150, 255; | |
| bra.uni $L__BB5_17; | |
| $L__BB5_10: | |
| setp.eq.s32 %p12, %r20, 0; | |
| @%p12 bra $L__BB5_12; | |
| add.s32 %r150, %r20, 112; | |
| mov.u32 %r151, %r148; | |
| bra.uni $L__BB5_17; | |
| $L__BB5_12: | |
| setp.eq.s32 %p13, %r21, 0; | |
| mov.u32 %r150, 0; | |
| mov.u32 %r151, %r150; | |
| @%p13 bra $L__BB5_17; | |
| mov.u32 %r150, 113; | |
| $L__BB5_14: | |
| shl.b32 %r26, %r148, 1; | |
| add.s32 %r150, %r150, -1; | |
| and.b32 %r63, %r148, 4194304; | |
| setp.eq.s32 %p14, %r63, 0; | |
| mov.u32 %r148, %r26; | |
| @%p14 bra $L__BB5_14; | |
| and.b32 %r151, %r26, 8388604; | |
| $L__BB5_17: | |
| shl.b32 %r65, %r150, 23; | |
| shl.b32 %r66, %r149, 31; | |
| or.b32 %r67, %r65, %r66; | |
| or.b32 %r68, %r67, %r151; | |
| mov.b32 %f142, %r68; | |
| mov.b32 %f141, %r18; | |
| $L__BB5_18: | |
| selp.f32 %f18, 0f3F800000, 0f00000000, %p1; | |
| abs.f32 %f19, %f141; | |
| setp.lt.f32 %p16, %f19, 0f7F800000; | |
| selp.f32 %f5, %f141, 0f00000000, %p16; | |
| setp.gt.f32 %p17, %f5, 0f00000000; | |
| setp.lt.f32 %p18, %f5, 0f00000000; | |
| selp.f32 %f20, 0fBF800000, 0f00000000, %p18; | |
| selp.f32 %f6, 0f3F800000, %f20, %p17; | |
| abs.f32 %f21, %f5; | |
| mov.b32 %r69, %f18; | |
| mov.u32 %r70, 2; | |
| mov.u32 %r71, 31; | |
| mov.u32 %r72, 16; | |
| mov.u32 %r73, -1; | |
| shfl.sync.down.b32 %r74|%p19, %r69, %r72, %r71, %r73; | |
| mov.b32 %f22, %r74; | |
| add.f32 %f23, %f18, %f22; | |
| mov.b32 %r75, %f6; | |
| shfl.sync.down.b32 %r76|%p20, %r75, %r72, %r71, %r73; | |
| mov.b32 %f24, %r76; | |
| add.f32 %f25, %f6, %f24; | |
| mov.b32 %r77, %f21; | |
| shfl.sync.down.b32 %r78|%p21, %r77, %r72, %r71, %r73; | |
| mov.b32 %f26, %r78; | |
| add.f32 %f27, %f21, %f26; | |
| mov.b32 %r79, %f23; | |
| mov.u32 %r80, 8; | |
| shfl.sync.down.b32 %r81|%p22, %r79, %r80, %r71, %r73; | |
| mov.b32 %f28, %r81; | |
| add.f32 %f29, %f23, %f28; | |
| mov.b32 %r82, %f25; | |
| shfl.sync.down.b32 %r83|%p23, %r82, %r80, %r71, %r73; | |
| mov.b32 %f30, %r83; | |
| add.f32 %f31, %f25, %f30; | |
| mov.b32 %r84, %f27; | |
| shfl.sync.down.b32 %r85|%p24, %r84, %r80, %r71, %r73; | |
| mov.b32 %f32, %r85; | |
| add.f32 %f33, %f27, %f32; | |
| mov.b32 %r86, %f29; | |
| mov.u32 %r87, 4; | |
| shfl.sync.down.b32 %r88|%p25, %r86, %r87, %r71, %r73; | |
| mov.b32 %f34, %r88; | |
| add.f32 %f35, %f29, %f34; | |
| mov.b32 %r89, %f31; | |
| shfl.sync.down.b32 %r90|%p26, %r89, %r87, %r71, %r73; | |
| mov.b32 %f36, %r90; | |
| add.f32 %f37, %f31, %f36; | |
| mov.b32 %r91, %f33; | |
| shfl.sync.down.b32 %r92|%p27, %r91, %r87, %r71, %r73; | |
| mov.b32 %f38, %r92; | |
| add.f32 %f39, %f33, %f38; | |
| mov.b32 %r93, %f35; | |
| shfl.sync.down.b32 %r94|%p28, %r93, %r70, %r71, %r73; | |
| mov.b32 %f40, %r94; | |
| add.f32 %f41, %f35, %f40; | |
| mov.b32 %r95, %f37; | |
| shfl.sync.down.b32 %r96|%p29, %r95, %r70, %r71, %r73; | |
| mov.b32 %f42, %r96; | |
| add.f32 %f43, %f37, %f42; | |
| mov.b32 %r97, %f39; | |
| shfl.sync.down.b32 %r98|%p30, %r97, %r70, %r71, %r73; | |
| mov.b32 %f44, %r98; | |
| add.f32 %f45, %f39, %f44; | |
| mov.b32 %r99, %f41; | |
| mov.u32 %r100, 1; | |
| shfl.sync.down.b32 %r101|%p31, %r99, %r100, %r71, %r73; | |
| mov.b32 %f46, %r101; | |
| add.f32 %f47, %f41, %f46; | |
| mov.b32 %r102, %f43; | |
| shfl.sync.down.b32 %r103|%p32, %r102, %r100, %r71, %r73; | |
| mov.b32 %f48, %r103; | |
| add.f32 %f49, %f43, %f48; | |
| mov.b32 %r104, %f45; | |
| shfl.sync.down.b32 %r105|%p33, %r104, %r100, %r71, %r73; | |
| mov.b32 %f50, %r105; | |
| add.f32 %f51, %f45, %f50; | |
| mov.b32 %r106, %f47; | |
| mov.u32 %r107, 0; | |
| shfl.sync.idx.b32 %r34|%p34, %r106, %r107, %r71, %r73; | |
| mov.b32 %r108, %f49; | |
| shfl.sync.idx.b32 %r35|%p35, %r108, %r107, %r71, %r73; | |
| mov.b32 %r109, %f51; | |
| shfl.sync.idx.b32 %r36|%p2, %r109, %r107, %r71, %r73; | |
| and.b32 %r110, %r1, 31; | |
| setp.ne.s32 %p36, %r110, 0; | |
| @%p36 bra $L__BB5_20; | |
| shr.u32 %r111, %r1, 3; | |
| and.b32 %r112, %r111, 536870908; | |
| mov.u32 %r113, _ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction; | |
| add.s32 %r114, %r113, %r112; | |
| st.shared.u32 [%r114], %r35; | |
| mov.u32 %r115, _ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active; | |
| add.s32 %r116, %r115, %r112; | |
| st.shared.u32 [%r116], %r34; | |
| $L__BB5_20: | |
| bar.sync 0; | |
| @%p3 bra $L__BB5_34; | |
| abs.f32 %f52, %f142; | |
| setp.lt.f32 %p38, %f52, 0f7F800000; | |
| selp.f32 %f7, %f142, 0f00000000, %p38; | |
| mov.f32 %f53, 0f00000000; | |
| mov.b32 %f54, %r34; | |
| mov.f32 %f55, 0f3F800000; | |
| max.f32 %f56, %f54, %f55; | |
| ld.shared.f32 %f57, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction]; | |
| ld.shared.f32 %f58, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active]; | |
| ld.shared.f32 %f59, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+4]; | |
| ld.shared.f32 %f60, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+4]; | |
| ld.shared.f32 %f61, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+8]; | |
| ld.shared.f32 %f62, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+8]; | |
| ld.shared.f32 %f63, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+12]; | |
| ld.shared.f32 %f64, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+12]; | |
| ld.shared.f32 %f65, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+16]; | |
| ld.shared.f32 %f66, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+16]; | |
| ld.shared.f32 %f67, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+20]; | |
| ld.shared.f32 %f68, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+20]; | |
| ld.shared.f32 %f69, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+24]; | |
| ld.shared.f32 %f70, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+24]; | |
| ld.shared.f32 %f71, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE17s_block_direction+28]; | |
| ld.shared.f32 %f72, [_ZZ17qpola_kernel_implI13__nv_fp8_e5m2EvPT_PKS1_ffiE14s_block_active+28]; | |
| mov.b32 %f73, %r35; | |
| div.rn.f32 %f74, %f73, %f56; | |
| mul.f32 %f75, %f6, %f74; | |
| sub.f32 %f76, %f55, %f75; | |
| max.f32 %f77, %f53, %f76; | |
| add.f32 %f78, %f58, 0f00000000; | |
| add.f32 %f79, %f78, %f60; | |
| add.f32 %f80, %f79, %f62; | |
| add.f32 %f81, %f80, %f64; | |
| add.f32 %f82, %f81, %f66; | |
| add.f32 %f83, %f82, %f68; | |
| add.f32 %f84, %f83, %f70; | |
| add.f32 %f85, %f84, %f72; | |
| max.f32 %f86, %f85, %f55; | |
| add.f32 %f87, %f57, 0f00000000; | |
| add.f32 %f88, %f87, %f59; | |
| add.f32 %f89, %f88, %f61; | |
| add.f32 %f90, %f89, %f63; | |
| add.f32 %f91, %f90, %f65; | |
| add.f32 %f92, %f91, %f67; | |
| add.f32 %f93, %f92, %f69; | |
| add.f32 %f94, %f93, %f71; | |
| div.rn.f32 %f95, %f94, %f86; | |
| mul.f32 %f96, %f6, %f95; | |
| sub.f32 %f97, %f55, %f96; | |
| max.f32 %f98, %f53, %f97; | |
| add.f32 %f8, %f77, %f98; | |
| mov.b32 %f99, %r36; | |
| div.rn.f32 %f100, %f99, %f56; | |
| add.f32 %f101, %f100, %f15; | |
| div.rn.f32 %f102, %f5, %f101; | |
| mul.f32 %f9, %f102, 0f3E800000; | |
| abs.f32 %f10, %f9; | |
| setp.ltu.f32 %p39, %f10, 0f3F19999A; | |
| @%p39 bra $L__BB5_23; | |
| bra.uni $L__BB5_22; | |
| $L__BB5_23: | |
| mul.f32 %f111, %f9, %f9; | |
| mov.f32 %f112, 0fBD563CAE; | |
| mov.f32 %f113, 0f3C80F082; | |
| fma.rn.f32 %f114, %f113, %f111, %f112; | |
| mov.f32 %f115, 0f3E085941; | |
| fma.rn.f32 %f116, %f114, %f111, %f115; | |
| mov.f32 %f117, 0fBEAAA9ED; | |
| fma.rn.f32 %f118, %f116, %f111, %f117; | |
| fma.rn.f32 %f120, %f118, %f111, %f53; | |
| fma.rn.f32 %f143, %f120, %f9, %f9; | |
| bra.uni $L__BB5_24; | |
| $L__BB5_22: | |
| mul.f32 %f103, %f10, 0f4038AA3B; | |
| ex2.approx.ftz.f32 %f104, %f103; | |
| add.f32 %f105, %f104, 0f3F800000; | |
| rcp.approx.ftz.f32 %f107, %f105; | |
| mov.f32 %f108, 0fC0000000; | |
| fma.rn.f32 %f109, %f107, %f108, %f55; | |
| setp.ge.f32 %p40, %f10, 0f41102CB4; | |
| selp.f32 %f110, 0f3F800000, %f109, %p40; | |
| mov.b32 %r117, %f110; | |
| mov.b32 %r118, %f9; | |
| and.b32 %r119, %r118, -2147483648; | |
| or.b32 %r120, %r119, %r117; | |
| mov.b32 %f143, %r120; | |
| $L__BB5_24: | |
| mul.f32 %f121, %f8, 0f3F000000; | |
| mul.f32 %f122, %f143, 0fC0800000; | |
| mul.f32 %f123, %f122, %f14; | |
| fma.rn.f32 %f124, %f121, 0fBEFD70A4, 0f3F800000; | |
| mov.f32 %f125, 0f3F800000; | |
| mov.f32 %f126, 0f3C23D70A; | |
| max.f32 %f127, %f124, %f126; | |
| min.f32 %f128, %f127, %f125; | |
| fma.rn.f32 %f129, %f128, %f123, %f7; | |
| abs.f32 %f130, %f129; | |
| setp.lt.f32 %p41, %f130, 0f7F800000; | |
| selp.f32 %f131, %f129, %f7, %p41; | |
| mov.f32 %f132, 0f47600000; | |
| min.f32 %f133, %f132, %f131; | |
| mov.f32 %f134, 0fC7600000; | |
| max.f32 %f135, %f134, %f133; | |
| fma.rn.f32 %f136, %f121, 0f3E4CCCCD, 0f3F800000; | |
| mul.f32 %f137, %f136, 0f3D800000; | |
| mad.lo.s32 %r122, %r2, 1664525, 7271263; | |
| and.b32 %r123, %r122, 16777215; | |
| cvt.rn.f32.u32 %f138, %r123; | |
| fma.rn.f32 %f139, %f138, 0f33800000, 0fBF000000; | |
| fma.rn.f32 %f140, %f139, %f137, %f135; | |
| mov.b32 %r124, %f140; | |
| and.b32 %r125, %r124, 2147483647; | |
| setp.gt.u32 %p42, %r125, 2139095040; | |
| cvt.f64.f32 %fd1, %f140; | |
| mov.b64 %rd13, %fd1; | |
| selp.b64 %rd2, 9223372036317904896, %rd13, %p42; | |
| shr.u64 %rd14, %rd2, 52; | |
| cvt.u32.u64 %r126, %rd14; | |
| add.s32 %r37, %r126, 16; | |
| shr.u64 %rd15, %rd2, 50; | |
| cvt.u32.u64 %r38, %rd15; | |
| and.b32 %r39, %r38, 3; | |
| and.b64 %rd3, %rd2, 9223372036854775807; | |
| setp.lt.u64 %p43, %rd3, 4530621225134718977; | |
| mov.u32 %r152, 0; | |
| @%p43 bra $L__BB5_33; | |
| setp.gt.u64 %p44, %rd3, 9218868437227405312; | |
| @%p44 bra $L__BB5_32; | |
| bra.uni $L__BB5_26; | |
| $L__BB5_32: | |
| or.b32 %r152, %r39, 126; | |
| bra.uni $L__BB5_33; | |
| $L__BB5_26: | |
| setp.gt.u64 %p45, %rd3, 4678677062884524031; | |
| mov.u32 %r152, 123; | |
| @%p45 bra $L__BB5_33; | |
| setp.lt.u64 %p46, %rd3, 4544132024016830464; | |
| @%p46 bra $L__BB5_29; | |
| bra.uni $L__BB5_28; | |
| $L__BB5_29: | |
| cvt.u16.u32 %rs9, %r37; | |
| mov.u16 %rs10, 1; | |
| sub.s16 %rs11, %rs10, %rs9; | |
| cvt.u32.u16 %r134, %rs11; | |
| and.b32 %r135, %r134, 255; | |
| or.b32 %r136, %r39, 4; | |
| shr.u32 %r152, %r136, %r135; | |
| mov.u64 %rd17, 1125899906842624; | |
| shl.b64 %rd18, %rd17, %r135; | |
| add.s64 %rd19, %rd18, -1; | |
| or.b64 %rd20, %rd2, 4503599627370496; | |
| and.b64 %rd4, %rd19, %rd20; | |
| mov.u64 %rd21, 562949953421312; | |
| shl.b64 %rd5, %rd21, %r135; | |
| setp.gt.u64 %p52, %rd4, %rd5; | |
| @%p52 bra $L__BB5_31; | |
| setp.ne.s64 %p53, %rd4, %rd5; | |
| and.b32 %r137, %r152, 1; | |
| setp.eq.b32 %p54, %r137, 1; | |
| not.pred %p55, %p54; | |
| or.pred %p56, %p55, %p53; | |
| @%p56 bra $L__BB5_33; | |
| $L__BB5_31: | |
| and.b32 %r138, %r152, 255; | |
| add.s32 %r152, %r138, 1; | |
| bra.uni $L__BB5_33; | |
| $L__BB5_28: | |
| shl.b32 %r128, %r37, 2; | |
| and.b32 %r129, %r128, 1020; | |
| or.b32 %r130, %r129, %r39; | |
| and.b64 %rd16, %rd2, 1125899906842623; | |
| setp.gt.u64 %p47, %rd16, 562949953421312; | |
| setp.eq.s64 %p48, %rd16, 562949953421312; | |
| and.b32 %r131, %r38, 1; | |
| setp.eq.b32 %p49, %r131, 1; | |
| and.pred %p50, %p48, %p49; | |
| or.pred %p51, %p47, %p50; | |
| and.b32 %r132, %r130, 255; | |
| add.s32 %r133, %r132, 1; | |
| selp.b32 %r152, %r133, %r130, %p51; | |
| $L__BB5_33: | |
| shr.u64 %rd22, %rd2, 63; | |
| cvt.u32.u64 %r139, %rd22; | |
| shl.b32 %r140, %r139, 7; | |
| or.b32 %r141, %r152, %r140; | |
| st.global.u8 [%rd1], %r141; | |
| $L__BB5_34: | |
| ret; | |
| } | |