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

.version 8.2
.target sm_89
.address_size 64

	// .globl	triton__0d1d2de
.extern .func __assertfail
(
	.param .b64 __assertfail_param_0,
	.param .b64 __assertfail_param_1,
	.param .b32 __assertfail_param_2,
	.param .b64 __assertfail_param_3,
	.param .b64 __assertfail_param_4
)
;
.global .align 1 .b8 assertFunc_0[25] = {95, 99, 97, 108, 108, 95, 119, 105, 116, 104, 95, 102, 114, 97, 109, 101, 115, 95, 114, 101, 109, 111, 118, 101, 100};
.global .align 1 .b8 assertFile_0[38] = {60, 102, 114, 111, 122, 101, 110, 32, 105, 109, 112, 111, 114, 116, 108, 105, 98, 46, 95, 98, 111, 111, 116, 115, 116, 114, 97, 112, 95, 101, 120, 116, 101, 114, 110, 97, 108, 62};
.global .align 1 .b8 assertMessage_0[38] = {105, 110, 100, 101, 120, 32, 111, 117, 116, 32, 111, 102, 32, 98, 111, 117, 110, 100, 115, 58, 32, 48, 32, 60, 61, 32, 116, 109, 112, 55, 32, 60, 32, 53, 48, 50, 53, 55};
.extern .shared .align 1 .b8 global_smem[];

.visible .entry triton__0d1d2de(
	.param .u64 triton__0d1d2de_param_0,
	.param .u64 triton__0d1d2de_param_1,
	.param .u64 triton__0d1d2de_param_2
)
.maxntid 128, 1, 1
{
	.reg .pred 	%p<24>;
	.reg .b16 	%rs<21>;
	.reg .b32 	%r<21>;
	.reg .b64 	%rd<58>;
	.loc	1 18 0
$L__func_begin0:
	.loc	1 18 0

	ld.param.u64 	%rd9, [triton__0d1d2de_param_1];
	ld.param.u64 	%rd16, [triton__0d1d2de_param_0];
$L__tmp0:
	.loc	1 21 36
	mov.u32 	%r4, %tid.x;
	and.b32  	%r1, %r4, 127;
	shl.b32 	%r2, %r1, 1;
	or.b32  	%r5, %r2, 1;
	or.b32  	%r6, %r2, 256;
	.loc	1 20 28
	mov.u32 %r3, %ctaid.x;
	.loc	1 20 46
	mul.wide.s32 	%rd1, %r3, 512;
	cvt.u64.u32 	%rd17, %r2;
	cvt.u64.u32 	%rd18, %r6;
	.loc	1 21 23
	or.b64  	%rd2, %rd1, %rd17;
	or.b64  	%rd3, %rd1, %rd18;
	.loc	1 24 30
	shl.b64 	%rd19, %rd2, 3;
	add.s64 	%rd12, %rd16, %rd19;
	add.s64 	%rd15, %rd12, 2048;
	mov.pred 	%p20, -1;
	.loc	1 24 35
	mov.u64 %rd10, 0x0;
	mov.u64 %rd11, 0x0;
	@%p20 ld.global.v2.b64 { %rd10, %rd11 }, [ %rd12 + 0 ];
	mov.u64 %rd13, 0x0;
	mov.u64 %rd14, 0x0;
	@%p20 ld.global.v2.b64 { %rd13, %rd14 }, [ %rd15 + 0 ];
	.loc	1 26 19
	setp.eq.s64 	%p3, %rd14, -1;
	setp.eq.s64 	%p4, %rd13, -1;
	setp.eq.s64 	%p5, %rd11, -1;
	setp.eq.s64 	%p6, %rd10, -1;
	.loc	1 28 32
	selp.b64 	%rd20, 0, %rd10, %p6;
	selp.b64 	%rd21, 0, %rd11, %p5;
	selp.b64 	%rd22, 0, %rd13, %p4;
	selp.b64 	%rd23, 0, %rd14, %p3;
	.loc	1 29 18
	add.s64 	%rd24, %rd23, 50257;
	add.s64 	%rd25, %rd22, 50257;
	add.s64 	%rd26, %rd21, 50257;
	add.s64 	%rd27, %rd20, 50257;
	.loc	1 30 18
	setp.lt.s64 	%p7, %rd23, 0;
	setp.lt.s64 	%p8, %rd22, 0;
	setp.lt.s64 	%p9, %rd21, 0;
	setp.lt.s64 	%p10, %rd20, 0;
	.loc	1 31 32
	selp.b64 	%rd7, %rd27, %rd20, %p10;
	selp.b64 	%rd6, %rd26, %rd21, %p9;
	selp.b64 	%rd5, %rd25, %rd22, %p8;
	selp.b64 	%rd4, %rd24, %rd23, %p7;
	.loc	1 32 36
	setp.lt.u64 	%p11, %rd4, 50257;
	setp.lt.u64 	%p12, %rd5, 50257;
	setp.lt.u64 	%p13, %rd6, 50257;
	setp.lt.u64 	%p14, %rd7, 50257;
	mov.u32 	%r7, global_smem;
	add.s32 	%r8, %r7, %r2;
	selp.u16 	%rs1, 1, 0, %p14;
	st.shared.u8 	[%r8], %rs1;
	cvt.u64.u32 	%rd8, %r5;
	selp.u16 	%rs2, 1, 0, %p13;
	st.shared.u8 	[%r8+1], %rs2;
	bar.sync 	0;
	add.s32 	%r9, %r7, %r1;
	ld.shared.u8 	%rs3, [%r9];
	ld.shared.u8 	%rs4, [%r9+128];
	bar.sync 	0;
	selp.u16 	%rs5, 1, 0, %p12;
	st.shared.u8 	[%r8], %rs5;
	selp.u16 	%rs6, 1, 0, %p11;
	st.shared.u8 	[%r8+1], %rs6;
	bar.sync 	0;
	ld.shared.u8 	%rs7, [%r9];
	ld.shared.u8 	%rs8, [%r9+128];
	setp.eq.s16 	%p15, %rs7, 0;
	selp.u16 	%rs9, 1, 0, %p15;
	shl.b16 	%rs10, %rs9, 2;
	setp.eq.s16 	%p16, %rs8, 0;
	selp.u16 	%rs11, -1, 0, %p16;
	shl.b16 	%rs12, %rs11, 3;
	or.b16  	%rs13, %rs12, %rs10;
	setp.eq.s16 	%p17, %rs4, 0;
	selp.u16 	%rs14, 1, 0, %p17;
	setp.eq.s16 	%p18, %rs3, 0;
	selp.u16 	%rs15, -1, 0, %p18;
	shl.b16 	%rs16, %rs15, 1;
	or.b16  	%rs17, %rs14, %rs16;
	and.b16  	%rs18, %rs17, 3;
	or.b16  	%rs19, %rs18, %rs13;
	.loc	1 32 51
	and.b16  	%rs20, %rs19, 15;
	setp.eq.s16 	%p19, %rs20, 0;
	@%p19 bra 	$L__BB0_2;
	mov.u64 	%rd28, assertMessage_0;
	cvta.global.u64 	%rd29, %rd28;
	mov.u64 	%rd30, assertFile_0;
	cvta.global.u64 	%rd31, %rd30;
	mov.u64 	%rd32, assertFunc_0;
	cvta.global.u64 	%rd33, %rd32;
	mov.b32 	%r10, 883;
	mov.u64 	%rd34, 1;
	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b64 param0;
	st.param.b64 	[param0+0], %rd29;
	.param .b64 param1;
	st.param.b64 	[param1+0], %rd31;
	.param .b32 param2;
	st.param.b32 	[param2+0], %r10;
	.param .b64 param3;
	st.param.b64 	[param3+0], %rd33;
	.param .b64 param4;
	st.param.b64 	[param4+0], %rd34;
	call.uni 
	__assertfail, 
	(
	param0, 
	param1, 
	param2, 
	param3, 
	param4
	);
	} // callseq 0
$L__BB0_2:
	.loc	1 21 36
	or.b32  	%r15, %r2, 257;
	cvt.u64.u32 	%rd39, %r15;
	.loc	1 21 23
	or.b64  	%rd40, %rd1, %rd39;
	or.b64  	%rd41, %rd1, %rd8;
	.loc	1 34 25
	shl.b64 	%rd42, %rd7, 2;
	add.s64 	%rd43, %rd9, %rd42;
	mul.lo.s64 	%rd44, %rd2, 201028;
	add.s64 	%rd45, %rd43, %rd44;
	shl.b64 	%rd46, %rd6, 2;
	add.s64 	%rd47, %rd9, %rd46;
	mul.lo.s64 	%rd48, %rd41, 201028;
	add.s64 	%rd49, %rd47, %rd48;
	shl.b64 	%rd50, %rd5, 2;
	add.s64 	%rd51, %rd9, %rd50;
	mul.lo.s64 	%rd52, %rd3, 201028;
	add.s64 	%rd53, %rd51, %rd52;
	shl.b64 	%rd54, %rd4, 2;
	add.s64 	%rd55, %rd9, %rd54;
	mul.lo.s64 	%rd56, %rd40, 201028;
	add.s64 	%rd57, %rd55, %rd56;
	.loc	1 34 51
	bar.sync 	0;
	shl.b32 	%r16, %r2, 3;
	add.s32 	%r18, %r7, %r16;
	st.shared.u64 	[%r18], %rd45;
	st.shared.u64 	[%r18+8], %rd49;
	bar.sync 	0;
	shl.b32 	%r19, %r1, 3;
	add.s32 	%r20, %r7, %r19;
	ld.shared.u64 	%rd35, [%r20];
	ld.shared.u64 	%rd36, [%r20+1024];
	bar.sync 	0;
	st.shared.u64 	[%r18], %rd53;
	st.shared.u64 	[%r18+8], %rd57;
	bar.sync 	0;
	ld.shared.u64 	%rd37, [%r20];
	ld.shared.u64 	%rd38, [%r20+1024];
	mov.b32 	%r11, -1082130432;
	@%p20 st.global.b32 [ %rd35 + 0 ], { %r11 };
	@%p20 st.global.b32 [ %rd36 + 0 ], { %r11 };
	@%p20 st.global.b32 [ %rd37 + 0 ], { %r11 };
	@%p20 st.global.b32 [ %rd38 + 0 ], { %r11 };
	.loc	1 34 4
	ret;
$L__tmp1:
$L__func_end0:

}
	.file	1 "/tmp/torchinductor_root/hl/chlrkgpvvbdizdz7sllquet2j7zhtes6meh6kenrqxov26mswvw7.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 11
.b8 63
.b8 12
.b8 0
.b8 0
.b8 0
	}
	.section	.debug_info
	{
.b32 176
.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 99
.b8 104
.b8 108
.b8 114
.b8 107
.b8 103
.b8 112
.b8 118
.b8 118
.b8 98
.b8 100
.b8 105
.b8 122
.b8 100
.b8 122
.b8 55
.b8 115
.b8 108
.b8 108
.b8 113
.b8 117
.b8 101
.b8 116
.b8 50
.b8 106
.b8 55
.b8 122
.b8 104
.b8 116
.b8 101
.b8 115
.b8 54
.b8 109
.b8 101
.b8 104
.b8 54
.b8 107
.b8 101
.b8 110
.b8 114
.b8 113
.b8 120
.b8 111
.b8 118
.b8 50
.b8 54
.b8 109
.b8 115
.b8 119
.b8 118
.b8 119
.b8 55
.b8 46
.b8 112
.b8 121
.b8 0
.b32 .debug_line
.b8 47
.b8 116
.b8 109
.b8 112
.b8 47
.b8 116
.b8 111
.b8 114
.b8 99
.b8 104
.b8 105
.b8 110
.b8 100
.b8 117
.b8 99
.b8 116
.b8 111
.b8 114
.b8 95
.b8 114
.b8 111
.b8 111
.b8 116
.b8 47
.b8 104
.b8 108
.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 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 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 101
.b8 0
.b8 1
.b8 18
.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 180
.b32 125
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 95
.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 180
.b32 0
$L__pubTypes_end0:
	}
	.section	.debug_loc	{	}