//
// Generated by LLVM NVPTX Back-End
//

.version 8.2
.target sm_90a
.address_size 64

	// .globl	triton_softmax_kernel_0d1d234

.visible .entry triton_softmax_kernel_0d1d234(
	.param .u64 triton_softmax_kernel_0d1d234_param_0,
	.param .u64 triton_softmax_kernel_0d1d234_param_1,
	.param .u32 triton_softmax_kernel_0d1d234_param_2,
	.param .u32 triton_softmax_kernel_0d1d234_param_3,
	.param .u32 triton_softmax_kernel_0d1d234_param_4
)
.maxntid 128, 1, 1
{
	.reg .pred 	%p<5>;
	.reg .b32 	%r<22>;
	.reg .f32 	%f<13>;
	.reg .b64 	%rd<10>;
	.loc	1 659 0
$L__func_begin0:
	.loc	1 659 0

	ld.param.u64 	%rd3, [triton_softmax_kernel_0d1d234_param_0];
	ld.param.u64 	%rd4, [triton_softmax_kernel_0d1d234_param_1];
$L__tmp0:
	.loc	1 664 31
	mov.u32 	%r8, %tid.x;
	and.b32  	%r9, %r8, 3;
	ld.param.u32 	%r10, [triton_softmax_kernel_0d1d234_param_2];
	ld.param.u32 	%r11, [triton_softmax_kernel_0d1d234_param_3];
	.loc	1 663 28
	mov.u32 %r1, %ctaid.x;
	.loc	1 667 36
	mul.lo.s32 	%r12, %r1, %r10;
	ld.param.u32 	%r13, [triton_softmax_kernel_0d1d234_param_4];
	.loc	1 667 26
	mul.wide.s32 	%rd5, %r12, 4;
	add.s64 	%rd6, %rd3, %rd5;
	.loc	1 668 27
	mul.wide.u32 	%rd7, %r9, 4;
	add.s64 	%rd1, %rd6, %rd7;
	.loc	1 669 47
	setp.lt.s32 	%p1, %r9, %r13;
	mov.b32 	%r3, -8388608;
	.loc	1 669 20
	mov.u32 %r2, 0x0;
	@%p1 ld.global.b32 { %r2 }, [ %rd1 + 0 ];
	@!%p1 mov.u32 %r2, %r3;
	mov.b32 	%f3, %r2;
$L__tmp1:
	.loc	2 169 40
	shfl.sync.bfly.b32	%r14, %r2, 2, 31, -1;
	mov.b32 	%f4, %r14;
$L__tmp2:
	.loc	2 124 23
	max.f32 	%f5, %f3, %f4;
$L__tmp3:
	.loc	2 169 40
	mov.b32 	%r15, %f5;
	shfl.sync.bfly.b32	%r16, %r15, 1, 31, -1;
	mov.b32 	%f6, %r16;
$L__tmp4:
	.loc	2 124 23
	max.f32 	%f7, %f5, %f6;
$L__tmp5:
	.loc	1 672 20
	sub.f32 	%f8, %f3, %f7;
	.loc	1 673 23
	mul.f32 	%f2, %f8, 0f3FB8AA3B;
	ex2.approx.f32 %f1, %f2;
$L__tmp6:
	.loc	2 243 36
	mov.b32 	%r5, %f1;
	shfl.sync.bfly.b32	%r17, %r5, 2, 31, -1;
	mov.b32 	%f9, %r17;
$L__tmp7:
	.loc	2 233 15
	add.f32 	%f10, %f1, %f9;
$L__tmp8:
	.loc	2 243 36
	mov.b32 	%r18, %f10;
	shfl.sync.bfly.b32	%r19, %r18, 1, 31, -1;
	mov.b32 	%f11, %r19;
$L__tmp9:
	.loc	2 233 15
	add.f32 	%f12, %f10, %f11;
$L__tmp10:
	.loc	1 675 24
	mov.b32 	%r6, %f12;
	div.full.f32 %r7, %r5, %r6;
	.loc	1 678 36
	mul.lo.s32 	%r20, %r1, %r11;
	.loc	1 678 26
	mul.wide.s32 	%rd8, %r20, 4;
	add.s64 	%rd9, %rd4, %rd8;
	.loc	1 679 27
	add.s64 	%rd2, %rd9, %rd7;
	.loc	1 680 21
	and.b32  	%r21, %r8, 124;
	setp.eq.s32 	%p4, %r21, 0;
	and.pred  	%p3, %p4, %p1;
	@%p3 st.global.b32 [ %rd2 + 0 ], { %r7 };
	.loc	1 680 4
	ret;
$L__tmp11:
$L__func_end0:

}
	.file	1 "/home/c-pliang/spring2024-lectures/lecture_06.py"
	.file	2 "/home/c-pliang/venv/lib/python3.10/site-packages/triton/language/standard.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 5
.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 5
.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 5
.b8 87
.b8 11
.b8 0
.b8 0
.b8 6
.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 0
	}
	.section	.debug_info
	{
.b32 328
.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 108
.b8 101
.b8 99
.b8 116
.b8 117
.b8 114
.b8 101
.b8 95
.b8 48
.b8 54
.b8 46
.b8 112
.b8 121
.b8 0
.b32 .debug_line
.b8 47
.b8 104
.b8 111
.b8 109
.b8 101
.b8 47
.b8 99
.b8 45
.b8 112
.b8 108
.b8 105
.b8 97
.b8 110
.b8 103
.b8 47
.b8 115
.b8 112
.b8 114
.b8 105
.b8 110
.b8 103
.b8 50
.b8 48
.b8 50
.b8 52
.b8 45
.b8 108
.b8 101
.b8 99
.b8 116
.b8 117
.b8 114
.b8 101
.b8 115
.b8 0
.b8 1
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 2
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.b8 115
.b8 111
.b8 102
.b8 116
.b8 109
.b8 97
.b8 120
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 51
.b8 52
.b8 0
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.b8 115
.b8 111
.b8 102
.b8 116
.b8 109
.b8 97
.b8 120
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 51
.b8 52
.b8 0
.b8 1
.b8 147
.b8 2
.b8 1
.b8 1
.b8 3
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b32 91
.b8 4
.b32 91
.b64 $L__tmp1
.b64 $L__tmp4
.b8 2
.b8 160
.b8 2
.b8 27
.b8 5
.b32 91
.b64 $L__tmp2
.b64 $L__tmp5
.b8 2
.b8 160
.b8 2
.b8 27
.b8 6
.b32 91
.b64 $L__tmp2
.b64 $L__tmp5
.b8 2
.b8 169
.b8 40
.b8 0
.b8 4
.b32 91
.b64 $L__tmp6
.b64 $L__tmp9
.b8 2
.b8 162
.b8 2
.b8 25
.b8 5
.b32 91
.b64 $L__tmp7
.b64 $L__tmp10
.b8 2
.b8 162
.b8 2
.b8 25
.b8 6
.b32 91
.b64 $L__tmp7
.b64 $L__tmp10
.b8 2
.b8 243
.b8 36
.b8 0
.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 332
.b32 91
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.b8 115
.b8 111
.b8 102
.b8 116
.b8 109
.b8 97
.b8 120
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 51
.b8 52
.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 332
.b32 0
$L__pubTypes_end0:
	}
	.section	.debug_loc	{	}