scan3 PTX from nanokitchen. Why is it numerically unstable? Is that the gate multiplication? Something else?

scan3.ptx

//
// Generated by LLVM NVPTX Back-End
// See scan3 in https://github.com/proger/nanokitchen/blob/main/triscan.py
//

.version 8.1
.target sm_86
.address_size 64

	// .globl	scan3_0d1d2d
.extern .shared .align 1 .b8 global_smem[];

.visible .entry scan3_0d1d2d(
	.param .u64 scan3_0d1d2d_param_0,
	.param .u64 scan3_0d1d2d_param_1,
	.param .u64 scan3_0d1d2d_param_2
)
.maxntid 128, 1, 1
{
	.reg .pred 	%p<13>;
	.reg .b32 	%r<67>;
	.reg .f32 	%f<48>;
	.reg .b64 	%rd<58>;
	.loc	1 94 0
$L__func_begin0:
	.loc	1 94 0

	ld.param.u64 	%rd5, [scan3_0d1d2d_param_0];
	ld.param.u64 	%rd6, [scan3_0d1d2d_param_1];
$L__tmp0:
	.loc	1 101 27
	mov.u32 	%r29, %tid.x;
	and.b32  	%r30, %r29, 31;
	ld.param.u64 	%rd7, [scan3_0d1d2d_param_2];
	bfe.u32 	%r31, %r29, 5, 2;
	and.b32  	%r32, %r29, 63;
	.loc	1 100 32
	mov.u32 	%r33, %nctaid.y;
	.loc	1 100 56
	mov.u32 	%r34, %ctaid.x;
	.loc	1 100 80
	mov.u32 	%r35, %ctaid.y;
	.loc	1 100 66
	mad.lo.s32 	%r36, %r33, %r34, %r35;
	.loc	1 101 58
	shl.b32 	%r37, %r36, 6;
	.loc	1 101 46
	or.b32  	%r38, %r37, %r32;
	.loc	1 103 31
	mul.wide.s32 	%rd8, %r38, 4;
	add.s64 	%rd1, %rd6, %rd8;
	mov.pred 	%p1, -1;
	.loc	1 103 22
	mov.u32 %r1, 0x0;
	@%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];
	.loc	1 104 29
	add.s64 	%rd2, %rd5, %rd8;
	.loc	1 104 21
	mov.u32 %r4, 0x0;
	@%p1 ld.global.b32 { %r4 }, [ %rd2 + 0 ];
$L__tmp1:
	.loc	1 79 46
	cvt.u64.u32 	%rd9, %r1;
	.loc	1 80 13
	shl.b64 	%rd10, %rd9, 32;
	.loc	1 81 46
	cvt.s64.s32 	%rd11, %r4;
	.loc	1 82 15
	or.b64  	%rd12, %rd10, %rd11;
$L__tmp2:
	.loc	1 107 40
	mov.b32 	%f1, %r4;
	shfl.sync.up.b32 %r3, %r4, 0x1, 0x0, 0xffffffff;
	mov.b32 	%f2, %r3;
$L__tmp3:
	.loc	1 82 15
	{ .reg .b32 tmp; mov.b64 {tmp, %r6}, %rd12; }
$L__tmp4:
	.loc	1 107 40
	shfl.sync.up.b32 %r5, %r6, 0x1, 0x0, 0xffffffff;
	mov.b32 	%f3, %r5;
	mov.b32 	%f4, %r6;
$L__tmp5:
	.loc	1 88 18
	fma.rn.f32 	%f5, %f2, %f4, %f3;
	.loc	1 89 13
	mul.f32 	%f6, %f2, %f1;
$L__tmp6:
	.loc	1 79 29
	mov.b32 	%r39, %f5;
	.loc	1 79 46
	cvt.u64.u32 	%rd13, %r39;
	.loc	1 80 13
	shl.b64 	%rd14, %rd13, 32;
	.loc	1 81 29
	mov.b32 	%r40, %f6;
	.loc	1 81 46
	cvt.s64.s32 	%rd15, %r40;
	.loc	1 82 15
	or.b64  	%rd16, %rd14, %rd15;
$L__tmp7:
	.loc	1 107 40
	setp.eq.s32 	%p5, %r30, 0;
	selp.b64 	%rd17, %rd12, %rd16, %p5;
	cvt.u32.u64 	%r8, %rd17;
	mov.b32 	%f7, %r8;
	shfl.sync.up.b32 %r7, %r8, 0x2, 0x0, 0xffffffff;
	mov.b32 	%f8, %r7;
	{ .reg .b32 tmp; mov.b64 {tmp, %r10}, %rd17; }
	shfl.sync.up.b32 %r9, %r10, 0x2, 0x0, 0xffffffff;
	mov.b32 	%f9, %r9;
	mov.b32 	%f10, %r10;
$L__tmp8:
	.loc	1 88 18
	fma.rn.f32 	%f11, %f8, %f10, %f9;
	.loc	1 89 13
	mul.f32 	%f12, %f8, %f7;
$L__tmp9:
	.loc	1 79 29
	mov.b32 	%r41, %f11;
	.loc	1 79 46
	cvt.u64.u32 	%rd18, %r41;
	.loc	1 80 13
	shl.b64 	%rd19, %rd18, 32;
	.loc	1 81 29
	mov.b32 	%r42, %f12;
	.loc	1 81 46
	cvt.s64.s32 	%rd20, %r42;
	.loc	1 82 15
	or.b64  	%rd21, %rd19, %rd20;
$L__tmp10:
	.loc	1 107 40
	setp.lt.u32 	%p6, %r30, 2;
	selp.b64 	%rd22, %rd17, %rd21, %p6;
	cvt.u32.u64 	%r12, %rd22;
	mov.b32 	%f13, %r12;
	shfl.sync.up.b32 %r11, %r12, 0x4, 0x0, 0xffffffff;
	mov.b32 	%f14, %r11;
	{ .reg .b32 tmp; mov.b64 {tmp, %r14}, %rd22; }
	shfl.sync.up.b32 %r13, %r14, 0x4, 0x0, 0xffffffff;
	mov.b32 	%f15, %r13;
	mov.b32 	%f16, %r14;
$L__tmp11:
	.loc	1 88 18
	fma.rn.f32 	%f17, %f14, %f16, %f15;
	.loc	1 89 13
	mul.f32 	%f18, %f14, %f13;
$L__tmp12:
	.loc	1 79 29
	mov.b32 	%r43, %f17;
	.loc	1 79 46
	cvt.u64.u32 	%rd23, %r43;
	.loc	1 80 13
	shl.b64 	%rd24, %rd23, 32;
	.loc	1 81 29
	mov.b32 	%r44, %f18;
	.loc	1 81 46
	cvt.s64.s32 	%rd25, %r44;
	.loc	1 82 15
	or.b64  	%rd26, %rd24, %rd25;
$L__tmp13:
	.loc	1 107 40
	setp.lt.u32 	%p7, %r30, 4;
	selp.b64 	%rd27, %rd22, %rd26, %p7;
	cvt.u32.u64 	%r16, %rd27;
	mov.b32 	%f19, %r16;
	shfl.sync.up.b32 %r15, %r16, 0x8, 0x0, 0xffffffff;
	mov.b32 	%f20, %r15;
	{ .reg .b32 tmp; mov.b64 {tmp, %r18}, %rd27; }
	shfl.sync.up.b32 %r17, %r18, 0x8, 0x0, 0xffffffff;
	mov.b32 	%f21, %r17;
	mov.b32 	%f22, %r18;
$L__tmp14:
	.loc	1 88 18
	fma.rn.f32 	%f23, %f20, %f22, %f21;
	.loc	1 89 13
	mul.f32 	%f24, %f20, %f19;
$L__tmp15:
	.loc	1 79 29
	mov.b32 	%r45, %f23;
	.loc	1 79 46
	cvt.u64.u32 	%rd28, %r45;
	.loc	1 80 13
	shl.b64 	%rd29, %rd28, 32;
	.loc	1 81 29
	mov.b32 	%r46, %f24;
	.loc	1 81 46
	cvt.s64.s32 	%rd30, %r46;
	.loc	1 82 15
	or.b64  	%rd31, %rd29, %rd30;
$L__tmp16:
	.loc	1 107 40
	setp.lt.u32 	%p8, %r30, 8;
	selp.b64 	%rd32, %rd27, %rd31, %p8;
	cvt.u32.u64 	%r20, %rd32;
	mov.b32 	%f25, %r20;
	shfl.sync.up.b32 %r19, %r20, 0x10, 0x0, 0xffffffff;
	mov.b32 	%f26, %r19;
	{ .reg .b32 tmp; mov.b64 {tmp, %r22}, %rd32; }
	shfl.sync.up.b32 %r21, %r22, 0x10, 0x0, 0xffffffff;
	mov.b32 	%f27, %r21;
	mov.b32 	%f28, %r22;
$L__tmp17:
	.loc	1 88 18
	fma.rn.f32 	%f29, %f26, %f28, %f27;
	.loc	1 89 13
	mul.f32 	%f30, %f26, %f25;
$L__tmp18:
	.loc	1 79 29
	mov.b32 	%r47, %f29;
	.loc	1 79 46
	cvt.u64.u32 	%rd33, %r47;
	.loc	1 80 13
	shl.b64 	%rd34, %rd33, 32;
	.loc	1 81 29
	mov.b32 	%r48, %f30;
	.loc	1 81 46
	cvt.s64.s32 	%rd35, %r48;
	.loc	1 82 15
	or.b64  	%rd36, %rd34, %rd35;
$L__tmp19:
	.loc	1 107 40
	setp.lt.u32 	%p9, %r30, 16;
	selp.b64 	%rd3, %rd32, %rd36, %p9;
	setp.eq.s32 	%p3, %r30, 31;
	shl.b32 	%r49, %r31, 3;
	mov.u32 	%r50, global_smem;
	add.s32 	%r23, %r50, %r49;
	@%p3 st.shared.b64 [ %r23 + 0 ], %rd3;
	bar.sync 	0;
	setp.eq.s32 	%p10, %r31, 0;
	ld.shared.u32 	%rd37, [global_smem+4];
	shl.b64 	%rd38, %rd37, 32;
	ld.shared.u32 	%rd39, [global_smem];
	or.b64  	%rd40, %rd38, %rd39;
	ld.shared.u32 	%rd41, [global_smem+12];
$L__tmp20:
	.loc	1 71 33
	cvt.u32.u64 	%r51, %rd39;
	.loc	1 71 46
	mov.b32 	%f31, %r51;
	.loc	1 72 0
	cvt.u32.u64 	%r52, %rd37;
	.loc	1 72 39
	mov.b32 	%f32, %r52;
$L__tmp21:
	.loc	1 71 46
	ld.shared.f32 	%f33, [global_smem+8];
	.loc	1 72 0
	cvt.u32.u64 	%r53, %rd41;
	.loc	1 72 39
	mov.b32 	%f34, %r53;
$L__tmp22:
	.loc	1 88 18
	fma.rn.f32 	%f35, %f32, %f33, %f34;
	.loc	1 89 13
	mul.f32 	%f36, %f31, %f33;
$L__tmp23:
	.loc	1 79 29
	mov.b32 	%r54, %f35;
	.loc	1 79 46
	cvt.u64.u32 	%rd42, %r54;
	.loc	1 80 13
	shl.b64 	%rd43, %rd42, 32;
	.loc	1 81 29
	mov.b32 	%r55, %f36;
	.loc	1 81 46
	cvt.s64.s32 	%rd44, %r55;
	.loc	1 82 15
	or.b64  	%rd45, %rd43, %rd44;
$L__tmp24:
	.loc	1 107 40
	setp.lt.u32 	%p11, %r31, 2;
	selp.b64 	%rd46, %rd40, %rd45, %p11;
	ld.shared.u32 	%rd47, [global_smem+20];
$L__tmp25:
	.loc	1 72 0
	{ .reg .b32 tmp; mov.b64 {tmp, %r56}, %rd45; }
	.loc	1 72 39
	mov.b32 	%f37, %r56;
$L__tmp26:
	.loc	1 71 46
	ld.shared.f32 	%f38, [global_smem+16];
	.loc	1 72 0
	cvt.u32.u64 	%r57, %rd47;
	.loc	1 72 39
	mov.b32 	%f39, %r57;
$L__tmp27:
	.loc	1 88 18
	fma.rn.f32 	%f40, %f38, %f37, %f39;
	.loc	1 89 13
	mul.f32 	%f41, %f36, %f38;
$L__tmp28:
	.loc	1 79 29
	mov.b32 	%r58, %f40;
	.loc	1 79 46
	cvt.u64.u32 	%rd48, %r58;
	.loc	1 80 13
	shl.b64 	%rd49, %rd48, 32;
	.loc	1 81 29
	mov.b32 	%r59, %f41;
	.loc	1 81 46
	cvt.s64.s32 	%rd50, %r59;
	.loc	1 82 15
	or.b64  	%rd51, %rd49, %rd50;
$L__tmp29:
	.loc	1 107 40
	setp.eq.s32 	%p12, %r31, 3;
	selp.b64 	%rd52, %rd51, %rd46, %p12;
$L__tmp30:
	.loc	1 71 33
	cvt.u32.u64 	%r60, %rd3;
	.loc	1 71 46
	mov.b32 	%f42, %r60;
	.loc	1 72 0
	{ .reg .b32 tmp; mov.b64 {tmp, %r61}, %rd3; }
	.loc	1 72 39
	mov.b32 	%f43, %r61;
$L__tmp31:
	.loc	1 71 33
	cvt.u32.u64 	%r62, %rd52;
	.loc	1 71 46
	mov.b32 	%f44, %r62;
	.loc	1 72 0
	{ .reg .b32 tmp; mov.b64 {tmp, %r63}, %rd52; }
	.loc	1 72 39
	mov.b32 	%f45, %r63;
$L__tmp32:
	.loc	1 88 18
	fma.rn.f32 	%f46, %f44, %f43, %f45;
	.loc	1 89 13
	mul.f32 	%f47, %f44, %f42;
$L__tmp33:
	.loc	1 79 29
	mov.b32 	%r64, %f46;
	.loc	1 79 46
	cvt.u64.u32 	%rd53, %r64;
	.loc	1 80 13
	shl.b64 	%rd54, %rd53, 32;
	.loc	1 81 29
	mov.b32 	%r65, %f47;
	.loc	1 81 46
	cvt.s64.s32 	%rd55, %r65;
	.loc	1 82 15
	or.b64  	%rd56, %rd54, %rd55;
$L__tmp34:
	.loc	1 107 40
	selp.b64 	%rd57, %rd3, %rd56, %p10;
$L__tmp35:
	.loc	1 72 0
	{ .reg .b32 tmp; mov.b64 {tmp, %r28}, %rd57; }
$L__tmp36:
	.loc	1 109 23
	add.s64 	%rd4, %rd7, %rd8;
	.loc	1 109 32
	and.b32  	%r66, %r29, 64;
	setp.eq.s32 	%p4, %r66, 0;
	@%p4 st.global.b32 [ %rd4 + 0 ], { %r28 };
	.loc	1 109 4
	ret;
$L__tmp37:
$L__func_end0:

}
	.file	1 "/home/proger/nanokitchen/triscan.py"
	.section	.debug_abbrev
	{
.b8 1
.b8 17
.b8 1
.b8 37
.b8 8
.b8 19
.b8 5
.b8 3
.b8 8
.b8 16
.b8 6
.b8 27
.b8 8
.b8 180
.b8 66
.b8 12
.b8 17
.b8 1
.b8 18
.b8 1
.b8 0
.b8 0
.b8 2
.b8 46
.b8 0
.b8 135
.b8 64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 63
.b8 12
.b8 32
.b8 11
.b8 0
.b8 0
.b8 3
.b8 46
.b8 1
.b8 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 49
.b8 19
.b8 0
.b8 0
.b8 4
.b8 29
.b8 0
.b8 49
.b8 19
.b8 17
.b8 1
.b8 18
.b8 1
.b8 88
.b8 11
.b8 89
.b8 11
.b8 87
.b8 11
.b8 0
.b8 0
.b8 5
.b8 29
.b8 1
.b8 49
.b8 19
.b8 17
.b8 1
.b8 18
.b8 1
.b8 88
.b8 11
.b8 89
.b8 11
.b8 87
.b8 11
.b8 0
.b8 0
.b8 0
	}
	.section	.debug_info
	{
.b32 299
.b8 2
.b8 0
.b32 .debug_abbrev
.b8 8
.b8 1
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 0
.b8 2
.b8 0
.b8 116
.b8 114
.b8 105
.b8 115
.b8 99
.b8 97
.b8 110
.b8 46
.b8 112
.b8 121
.b8 0
.b32 .debug_line
.b8 47
.b8 104
.b8 111
.b8 109
.b8 101
.b8 47
.b8 112
.b8 114
.b8 111
.b8 103
.b8 101
.b8 114
.b8 47
.b8 110
.b8 97
.b8 110
.b8 111
.b8 107
.b8 105
.b8 116
.b8 99
.b8 104
.b8 101
.b8 110
.b8 0
.b8 1
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 2
.b8 115
.b8 99
.b8 97
.b8 110
.b8 51
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 0
.b8 115
.b8 99
.b8 97
.b8 110
.b8 51
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 0
.b8 1
.b8 94
.b8 1
.b8 1
.b8 3
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b32 78
.b8 4
.b32 78
.b64 $L__tmp1
.b64 $L__tmp4
.b8 1
.b8 106
.b8 29
.b8 4
.b32 78
.b64 $L__tmp5
.b64 $L__tmp33
.b8 1
.b8 107
.b8 40
.b8 5
.b32 78
.b64 $L__tmp6
.b64 $L__tmp34
.b8 1
.b8 107
.b8 40
.b8 4
.b32 78
.b64 $L__tmp6
.b64 $L__tmp34
.b8 1
.b8 90
.b8 21
.b8 4
.b32 78
.b64 $L__tmp20
.b64 $L__tmp31
.b8 1
.b8 86
.b8 22
.b8 4
.b32 78
.b64 $L__tmp21
.b64 $L__tmp32
.b8 1
.b8 87
.b8 22
.b8 0
.b8 4
.b32 78
.b64 $L__tmp35
.b64 $L__tmp36
.b8 1
.b8 108
.b8 43
.b8 0
.b8 0
	}
	.section	.debug_pubnames
	{
.b32 $L__pubNames_end0-$L__pubNames_start0
$L__pubNames_start0:
.b8 2
.b8 0
.b32 .debug_info
.b32 303
.b32 78
.b8 115
.b8 99
.b8 97
.b8 110
.b8 51
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 0
.b32 0
$L__pubNames_end0:
	}
	.section	.debug_pubtypes
	{
.b32 $L__pubTypes_end0-$L__pubTypes_start0
$L__pubTypes_start0:
.b8 2
.b8 0
.b32 .debug_info
.b32 303
.b32 0
$L__pubTypes_end0:
	}
	.section	.debug_loc	{	}