//
// 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;

}