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

.version 8.2
.target sm_90a
.address_size 64

	// .globl	triton_gelu_kernel_0d1d2de

.visible .entry triton_gelu_kernel_0d1d2de(
	.param .u64 triton_gelu_kernel_0d1d2de_param_0,
	.param .u64 triton_gelu_kernel_0d1d2de_param_1,
	.param .u32 triton_gelu_kernel_0d1d2de_param_2
)
.maxntid 128, 1, 1
{
	.reg .pred 	%p<5>;
	.reg .b32 	%r<49>;
	.reg .f32 	%f<113>;
	.reg .b64 	%rd<8>;
	.loc	1 484 0
$L__func_begin0:
	.loc	1 484 0

	ld.param.u64 	%rd5, [triton_gelu_kernel_0d1d2de_param_0];
	ld.param.u64 	%rd6, [triton_gelu_kernel_0d1d2de_param_1];
$L__tmp0:
	.loc	1 493 41
	mov.u32 	%r42, %tid.x;
	shl.b32 	%r43, %r42, 2;
	ld.param.u32 	%r44, [triton_gelu_kernel_0d1d2de_param_2];
	and.b32  	%r45, %r43, 508;
	.loc	1 489 24
	mov.u32 %r1, %ctaid.x;
	.loc	1 490 24
	shl.b32 	%r46, %r1, 10;
	.loc	1 493 28
	or.b32  	%r47, %r46, %r45;
	or.b32  	%r48, %r47, 512;
	.loc	1 496 21
	setp.lt.s32 	%p1, %r47, %r44;
	setp.lt.s32 	%p2, %r48, %r44;
	.loc	1 499 24
	mul.wide.s32 	%rd7, %r47, 4;
	add.s64 	%rd1, %rd5, %rd7;
	add.s64 	%rd2, %rd1, 2048;
	.loc	1 499 16
	mov.u32 %r2, 0x0;
	mov.u32 %r3, 0x0;
	mov.u32 %r4, 0x0;
	mov.u32 %r5, 0x0;
	@%p1 ld.global.v4.b32 { %r2, %r3, %r4, %r5 }, [ %rd1 + 0 ];
	mov.b32 	%f17, %r2;
	mov.b32 	%f18, %r3;
	mov.b32 	%f19, %r4;
	mov.b32 	%f20, %r5;
	mov.u32 %r6, 0x0;
	mov.u32 %r7, 0x0;
	mov.u32 %r8, 0x0;
	mov.u32 %r9, 0x0;
	@%p2 ld.global.v4.b32 { %r6, %r7, %r8, %r9 }, [ %rd2 + 0 ];
	mov.b32 	%f21, %r6;
	mov.b32 	%f22, %r7;
	mov.b32 	%f23, %r8;
	mov.b32 	%f24, %r9;
	.loc	1 502 37
	mul.f32 	%f25, %f17, 0f3D372713;
	mul.f32 	%f26, %f18, 0f3D372713;
	mul.f32 	%f27, %f19, 0f3D372713;
	mul.f32 	%f28, %f20, 0f3D372713;
	mul.f32 	%f29, %f21, 0f3D372713;
	mul.f32 	%f30, %f22, 0f3D372713;
	mul.f32 	%f31, %f23, 0f3D372713;
	mul.f32 	%f32, %f24, 0f3D372713;
	.loc	1 502 41
	mul.f32 	%f33, %f25, %f17;
	mul.f32 	%f34, %f26, %f18;
	mul.f32 	%f35, %f27, %f19;
	mul.f32 	%f36, %f28, %f20;
	mul.f32 	%f37, %f29, %f21;
	mul.f32 	%f38, %f30, %f22;
	mul.f32 	%f39, %f31, %f23;
	mul.f32 	%f40, %f32, %f24;
	.loc	1 502 26
	fma.rn.f32 	%f41, %f33, %f17, %f17;
	fma.rn.f32 	%f42, %f34, %f18, %f18;
	fma.rn.f32 	%f43, %f35, %f19, %f19;
	fma.rn.f32 	%f44, %f36, %f20, %f20;
	fma.rn.f32 	%f45, %f37, %f21, %f21;
	fma.rn.f32 	%f46, %f38, %f22, %f22;
	fma.rn.f32 	%f47, %f39, %f23, %f23;
	fma.rn.f32 	%f48, %f40, %f24, %f24;
	.loc	1 502 22
	mul.f32 	%f49, %f41, 0f3F4C422A;
	mul.f32 	%f50, %f42, 0f3F4C422A;
	mul.f32 	%f51, %f43, 0f3F4C422A;
	mul.f32 	%f52, %f44, 0f3F4C422A;
	mul.f32 	%f53, %f45, 0f3F4C422A;
	mul.f32 	%f54, %f46, 0f3F4C422A;
	mul.f32 	%f55, %f47, 0f3F4C422A;
	mul.f32 	%f56, %f48, 0f3F4C422A;
	.loc	1 503 21
	fma.rn.f32 	%f57, %f41, 0f3F4C422A, %f49;
	fma.rn.f32 	%f58, %f42, 0f3F4C422A, %f50;
	fma.rn.f32 	%f59, %f43, 0f3F4C422A, %f51;
	fma.rn.f32 	%f60, %f44, 0f3F4C422A, %f52;
	fma.rn.f32 	%f61, %f45, 0f3F4C422A, %f53;
	fma.rn.f32 	%f62, %f46, 0f3F4C422A, %f54;
	fma.rn.f32 	%f63, %f47, 0f3F4C422A, %f55;
	fma.rn.f32 	%f64, %f48, 0f3F4C422A, %f56;
	.loc	1 503 17
	mul.f32 	%f2, %f57, 0f3FB8AA3B;
	ex2.approx.f32 %f1, %f2;
	mul.f32 	%f4, %f58, 0f3FB8AA3B;
	ex2.approx.f32 %f3, %f4;
	mul.f32 	%f6, %f59, 0f3FB8AA3B;
	ex2.approx.f32 %f5, %f6;
	mul.f32 	%f8, %f60, 0f3FB8AA3B;
	ex2.approx.f32 %f7, %f8;
	mul.f32 	%f10, %f61, 0f3FB8AA3B;
	ex2.approx.f32 %f9, %f10;
	mul.f32 	%f12, %f62, 0f3FB8AA3B;
	ex2.approx.f32 %f11, %f12;
	mul.f32 	%f14, %f63, 0f3FB8AA3B;
	ex2.approx.f32 %f13, %f14;
	mul.f32 	%f16, %f64, 0f3FB8AA3B;
	ex2.approx.f32 %f15, %f16;
	.loc	1 504 18
	add.f32 	%f65, %f1, 0fBF800000;
	add.f32 	%f66, %f3, 0fBF800000;
	add.f32 	%f67, %f5, 0fBF800000;
	add.f32 	%f68, %f7, 0fBF800000;
	add.f32 	%f69, %f9, 0fBF800000;
	add.f32 	%f70, %f11, 0fBF800000;
	add.f32 	%f71, %f13, 0fBF800000;
	add.f32 	%f72, %f15, 0fBF800000;
	.loc	1 504 30
	add.f32 	%f73, %f1, 0f3F800000;
	add.f32 	%f74, %f3, 0f3F800000;
	add.f32 	%f75, %f5, 0f3F800000;
	add.f32 	%f76, %f7, 0f3F800000;
	add.f32 	%f77, %f9, 0f3F800000;
	add.f32 	%f78, %f11, 0f3F800000;
	add.f32 	%f79, %f13, 0f3F800000;
	add.f32 	%f80, %f15, 0f3F800000;
	.loc	1 504 24
	mov.b32 	%r11, %f65;
	mov.b32 	%r12, %f73;
	div.full.f32 %r10, %r11, %r12;
	mov.b32 	%f81, %r10;
	mov.b32 	%r14, %f66;
	mov.b32 	%r15, %f74;
	div.full.f32 %r13, %r14, %r15;
	mov.b32 	%f82, %r13;
	mov.b32 	%r17, %f67;
	mov.b32 	%r18, %f75;
	div.full.f32 %r16, %r17, %r18;
	mov.b32 	%f83, %r16;
	mov.b32 	%r20, %f68;
	mov.b32 	%r21, %f76;
	div.full.f32 %r19, %r20, %r21;
	mov.b32 	%f84, %r19;
	mov.b32 	%r23, %f69;
	mov.b32 	%r24, %f77;
	div.full.f32 %r22, %r23, %r24;
	mov.b32 	%f85, %r22;
	mov.b32 	%r26, %f70;
	mov.b32 	%r27, %f78;
	div.full.f32 %r25, %r26, %r27;
	mov.b32 	%f86, %r25;
	mov.b32 	%r29, %f71;
	mov.b32 	%r30, %f79;
	div.full.f32 %r28, %r29, %r30;
	mov.b32 	%f87, %r28;
	mov.b32 	%r32, %f72;
	mov.b32 	%r33, %f80;
	div.full.f32 %r31, %r32, %r33;
	mov.b32 	%f88, %r31;
	.loc	1 505 14
	mul.f32 	%f89, %f17, 0f3F000000;
	mul.f32 	%f90, %f18, 0f3F000000;
	mul.f32 	%f91, %f19, 0f3F000000;
	mul.f32 	%f92, %f20, 0f3F000000;
	mul.f32 	%f93, %f21, 0f3F000000;
	mul.f32 	%f94, %f22, 0f3F000000;
	mul.f32 	%f95, %f23, 0f3F000000;
	mul.f32 	%f96, %f24, 0f3F000000;
	.loc	1 505 23
	add.f32 	%f97, %f81, 0f3F800000;
	add.f32 	%f98, %f82, 0f3F800000;
	add.f32 	%f99, %f83, 0f3F800000;
	add.f32 	%f100, %f84, 0f3F800000;
	add.f32 	%f101, %f85, 0f3F800000;
	add.f32 	%f102, %f86, 0f3F800000;
	add.f32 	%f103, %f87, 0f3F800000;
	add.f32 	%f104, %f88, 0f3F800000;
	.loc	1 505 19
	mul.f32 	%f105, %f89, %f97;
	mul.f32 	%f106, %f90, %f98;
	mul.f32 	%f107, %f91, %f99;
	mul.f32 	%f108, %f92, %f100;
	mul.f32 	%f109, %f93, %f101;
	mul.f32 	%f110, %f94, %f102;
	mul.f32 	%f111, %f95, %f103;
	mul.f32 	%f112, %f96, %f104;
	.loc	1 508 21
	add.s64 	%rd3, %rd6, %rd7;
	add.s64 	%rd4, %rd3, 2048;
	.loc	1 508 30
	mov.b32 	%r34, %f105;
	mov.b32 	%r35, %f106;
	mov.b32 	%r36, %f107;
	mov.b32 	%r37, %f108;
	@%p1 st.global.v4.b32 [ %rd3 + 0 ], { %r34, %r35, %r36, %r37 };
	mov.b32 	%r38, %f109;
	mov.b32 	%r39, %f110;
	mov.b32 	%r40, %f111;
	mov.b32 	%r41, %f112;
	@%p2 st.global.v4.b32 [ %rd4 + 0 ], { %r38, %r39, %r40, %r41 };
	.loc	1 508 4
	ret;
$L__tmp1:
$L__func_end0:

}
	.file	1 "/home/c-pliang/spring2024-lectures/lecture_06.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 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 135
.b8 64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 5
.b8 63
.b8 12
.b8 0
.b8 0
.b8 0
	}
	.section	.debug_info
	{
.b32 165
.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
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.b8 103
.b8 101
.b8 108
.b8 117
.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 100
.b8 101
.b8 0
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.b8 103
.b8 101
.b8 108
.b8 117
.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 100
.b8 101
.b8 0
.b8 1
.b8 228
.b8 1
.b8 1
.b8 0
	}
	.section	.debug_pubnames
	{
.b32 $L__pubNames_end0-$L__pubNames_start0
$L__pubNames_start0:
.b8 2
.b8 0
.b32 .debug_info
.b32 169
.b32 91
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.b8 103
.b8 101
.b8 108
.b8 117
.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 100
.b8 101
.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 169
.b32 0
$L__pubTypes_end0:
	}
	.section	.debug_loc	{	}