Add files using upload-large-folder tool

.triton/dump/0471aff594c8c8b8715b81c529738739/triton_.ptx

@@ -0,0 +1,951 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2d3d4d5d6de7de
+.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_1[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_1[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_1[39] = {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, 49, 54, 32, 60, 32, 53, 48, 50, 53, 55};
+.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, 51, 32, 60, 32, 53, 48, 50, 53, 55};
+.extern .shared .align 1 .b8 global_smem[];
+.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};
+
+.visible .entry triton__0d1d2d3d4d5d6de7de(
+	.param .u64 triton__0d1d2d3d4d5d6de7de_param_0,
+	.param .u64 triton__0d1d2d3d4d5d6de7de_param_1,
+	.param .u64 triton__0d1d2d3d4d5d6de7de_param_2,
+	.param .u64 triton__0d1d2d3d4d5d6de7de_param_3,
+	.param .u64 triton__0d1d2d3d4d5d6de7de_param_4,
+	.param .u64 triton__0d1d2d3d4d5d6de7de_param_5,
+	.param .u32 triton__0d1d2d3d4d5d6de7de_param_6,
+	.param .u32 triton__0d1d2d3d4d5d6de7de_param_7
+)
+.maxntid 128, 1, 1
+{
+	.reg .pred 	%p<48>;
+	.reg .b16 	%rs<13>;
+	.reg .b32 	%r<158>;
+	.reg .f32 	%f<164>;
+	.reg .b64 	%rd<73>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd19, [triton__0d1d2d3d4d5d6de7de_param_5];
+	ld.param.u64 	%rd18, [triton__0d1d2d3d4d5d6de7de_param_4];
+	ld.param.u64 	%rd17, [triton__0d1d2d3d4d5d6de7de_param_3];
+	ld.param.u64 	%rd30, [triton__0d1d2d3d4d5d6de7de_param_0];
+	ld.param.u64 	%rd31, [triton__0d1d2d3d4d5d6de7de_param_1];
+$L__tmp0:
+	.loc	1 22 44
+	mov.u32 	%r13, %tid.x;
+	and.b32  	%r1, %r13, 31;
+	ld.param.u64 	%rd32, [triton__0d1d2d3d4d5d6de7de_param_2];
+	bfe.u32 	%r14, %r13, 5, 2;
+	bfe.u32 	%r15, %r13, 1, 4;
+	shl.b32 	%r16, %r14, 4;
+	or.b32  	%r2, %r16, %r15;
+	and.b32  	%r17, %r13, 63;
+	.loc	1 24 33
+	shl.b32 	%r18, %r13, 2;
+	and.b32  	%r3, %r18, 4;
+	and.b32  	%r4, %r13, 7;
+	.loc	1 31 36
+	shl.b32 	%r5, %r14, 2;
+	.loc	1 21 28
+	mov.u32 %r11, %ctaid.x;
+	.loc	1 21 33
+	shl.b32 	%r19, %r11, 6;
+	.loc	1 22 23
+	or.b32  	%r20, %r19, %r2;
+	or.b32  	%r21, %r19, %r17;
+	.loc	1 26 30
+	mul.wide.s32 	%rd33, %r20, 8;
+	add.s64 	%rd21, %rd30, %rd33;
+	mul.wide.s32 	%rd34, %r21, 8;
+	add.s64 	%rd29, %rd30, %rd34;
+	mov.pred 	%p1, -1;
+	.loc	1 26 35
+	mov.u64 %rd20, 0x0;
+	@%p1 ld.global.L1::evict_last.b64 { %rd20 }, [ %rd21 + 0 ];
+	mov.u64 %rd22, 0x0;
+	@%p1 ld.global.L1::evict_last.b64 { %rd22 }, [ %rd21 + 0 ];
+	mov.u64 %rd24, 0x0;
+	@%p1 ld.global.L1::evict_last.b64 { %rd24 }, [ %rd21 + 0 ];
+	mov.u64 %rd26, 0x0;
+	@%p1 ld.global.L1::evict_last.b64 { %rd26 }, [ %rd21 + 0 ];
+	mov.u64 %rd28, 0x0;
+	@%p1 ld.global.L1::evict_last.b64 { %rd28 }, [ %rd29 + 0 ];
+	.loc	1 27 18
+	bfe.s32 	%r22, %r11, 25, 1;
+	shr.u32 	%r23, %r22, 23;
+	add.s32 	%r24, %r20, %r23;
+	and.b32  	%r25, %r24, 16776704;
+	sub.s32 	%r26, %r20, %r25;
+	.loc	1 35 44
+	shl.b32 	%r27, %r26, 8;
+	.loc	1 37 22
+	add.s64 	%rd35, %rd28, 50257;
+	.loc	1 38 22
+	setp.lt.s64 	%p6, %rd20, 0;
+	setp.lt.s64 	%p7, %rd28, 0;
+	.loc	1 39 36
+	selp.b64 	%rd1, %rd35, %rd28, %p7;
+	.loc	1 41 44
+	shl.b64 	%rd36, %rd20, 8;
+	add.s64 	%rd37, %rd36, 12865792;
+	selp.b64 	%rd38, %rd37, %rd36, %p6;
+	.loc	1 31 36
+	and.b32  	%r28, %r13, 1;
+	mul.wide.u32 	%rd2, %r28, 16;
+	shl.b64 	%rd39, %rd38, 2;
+	or.b64  	%rd40, %rd2, %rd39;
+	add.s64 	%rd72, %rd31, %rd40;
+	shl.b32 	%r29, %r11, 14;
+	shl.b32 	%r30, %r14, 12;
+	or.b32  	%r31, %r29, %r30;
+	shl.b32 	%r32, %r15, 8;
+	or.b32  	%r33, %r31, %r32;
+	or.b32  	%r6, %r33, %r3;
+	or.b32  	%r34, %r27, %r3;
+	mul.wide.s32 	%rd41, %r34, 4;
+	add.s64 	%rd70, %rd32, %rd41;
+	mov.f32 	%f148, 0f00000000;
+	mov.b32 	%r156, -8;
+	mov.u64 	%rd68, %rd70;
+	mov.u64 	%rd69, %rd72;
+	mov.f32 	%f149, %f148;
+	mov.f32 	%f150, %f148;
+	mov.f32 	%f151, %f148;
+	mov.f32 	%f152, %f148;
+	mov.f32 	%f153, %f148;
+	mov.f32 	%f154, %f148;
+	mov.f32 	%f155, %f148;
+	mov.f32 	%f156, %f148;
+	mov.f32 	%f157, %f148;
+	mov.f32 	%f158, %f148;
+	mov.f32 	%f159, %f148;
+	mov.f32 	%f160, %f148;
+	mov.f32 	%f161, %f148;
+	mov.f32 	%f162, %f148;
+	mov.f32 	%f163, %f148;
+	bra.uni 	$L__BB0_1;
+$L__BB0_3:
+	.loc	1 0 0
+	mov.b32 	%f17, %r35;
+	mov.b32 	%f18, %r36;
+	mov.b32 	%f19, %r37;
+	mov.b32 	%f20, %r38;
+	cvt.u16.u32 	%rs1, %r43;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r43; }
+	cvt.u16.u32 	%rs3, %r44;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r44; }
+	cvt.f32.bf16 %r47, %rs1;
+	mov.b32 	%f21, %r47;
+	cvt.f32.bf16 %r48, %rs2;
+	mov.b32 	%f22, %r48;
+	cvt.f32.bf16 %r49, %rs3;
+	mov.b32 	%f23, %r49;
+	cvt.f32.bf16 %r50, %rs4;
+	mov.b32 	%f24, %r50;
+	.loc	1 41 52
+	mov.u32 %r54, 0x0;
+	mov.u32 %r55, 0x0;
+	mov.u32 %r56, 0x0;
+	mov.u32 %r57, 0x0;
+	@%p1 ld.global.L1::evict_last.v4.b32 { %r54, %r55, %r56, %r57 }, [ %rd69 + 0 ];
+	@!%p1 mov.u32 %r54, %r143;
+	@!%p1 mov.u32 %r55, %r143;
+	@!%p1 mov.u32 %r56, %r143;
+	@!%p1 mov.u32 %r57, %r143;
+	mov.b32 	%f56, %r54;
+	mov.b32 	%f57, %r55;
+	mov.b32 	%f58, %r56;
+	mov.b32 	%f59, %r57;
+	.loc	1 42 22
+	add.f32 	%f60, %f17, %f56;
+	add.f32 	%f61, %f18, %f57;
+	add.f32 	%f62, %f19, %f58;
+	add.f32 	%f63, %f20, %f59;
+	.loc	1 44 22
+	add.f32 	%f64, %f21, %f60;
+	add.f32 	%f65, %f22, %f61;
+	add.f32 	%f66, %f23, %f62;
+	add.f32 	%f67, %f24, %f63;
+$L__tmp1:
+	.loc	2 96 20
+	sub.f32 	%f68, %f64, %f160;
+	sub.f32 	%f69, %f65, %f161;
+	sub.f32 	%f70, %f66, %f162;
+	sub.f32 	%f71, %f67, %f163;
+	.loc	2 97 26
+	add.f32 	%f148, %f148, 0f3F800000;
+	add.f32 	%f149, %f149, 0f3F800000;
+	add.f32 	%f150, %f150, 0f3F800000;
+	add.f32 	%f151, %f151, 0f3F800000;
+	add.f32 	%f152, %f152, 0f3F800000;
+	add.f32 	%f153, %f153, 0f3F800000;
+	add.f32 	%f154, %f154, 0f3F800000;
+	add.f32 	%f155, %f155, 0f3F800000;
+	.loc	2 98 30
+	mov.b32 	%r63, %f68;
+	mov.b32 	%r64, %f148;
+	div.full.f32 %r62, %r63, %r64;
+	mov.b32 	%f72, %r62;
+	mov.b32 	%r66, %f69;
+	mov.b32 	%r67, %f149;
+	div.full.f32 %r65, %r66, %r67;
+	mov.b32 	%f73, %r65;
+	mov.b32 	%r69, %f70;
+	mov.b32 	%r70, %f150;
+	div.full.f32 %r68, %r69, %r70;
+	mov.b32 	%f74, %r68;
+	mov.b32 	%r72, %f71;
+	mov.b32 	%r73, %f151;
+	div.full.f32 %r71, %r72, %r73;
+	mov.b32 	%f75, %r71;
+	.loc	2 98 22
+	add.f32 	%f160, %f160, %f72;
+	add.f32 	%f161, %f161, %f73;
+	add.f32 	%f162, %f162, %f74;
+	add.f32 	%f163, %f163, %f75;
+	.loc	2 101 30
+	sub.f32 	%f76, %f64, %f160;
+	sub.f32 	%f77, %f65, %f161;
+	sub.f32 	%f78, %f66, %f162;
+	sub.f32 	%f79, %f67, %f163;
+$L__tmp2:
+	.loc	1 50 50
+	fma.rn.f32 	%f156, %f68, %f76, %f156;
+	fma.rn.f32 	%f157, %f69, %f77, %f157;
+	fma.rn.f32 	%f158, %f70, %f78, %f158;
+	fma.rn.f32 	%f159, %f71, %f79, %f159;
+	.loc	1 31 36
+	add.s32 	%r156, %r156, 8;
+	add.s64 	%rd69, %rd69, 32;
+	add.s64 	%rd68, %rd68, 32;
+	setp.lt.u32 	%p22, %r156, 248;
+	@%p22 bra 	$L__BB0_1;
+	bra.uni 	$L__BB0_4;
+$L__BB0_1:
+	.loc	1 40 40
+	setp.lt.u64 	%p16, %rd1, 50257;
+	mov.b32 	%r143, 0;
+	.loc	1 35 50
+	mov.u32 %r35, 0x0;
+	mov.u32 %r36, 0x0;
+	mov.u32 %r37, 0x0;
+	mov.u32 %r38, 0x0;
+	@%p1 ld.global.L1::evict_last.v4.b32 { %r35, %r36, %r37, %r38 }, [ %rd68 + 0 ];
+	@!%p1 mov.u32 %r35, %r143;
+	@!%p1 mov.u32 %r36, %r143;
+	@!%p1 mov.u32 %r37, %r143;
+	@!%p1 mov.u32 %r38, %r143;
+	.loc	1 36 34
+	add.s32 	%r51, %r6, %r156;
+	add.s32 	%r52, %r51, 8;
+	mul.wide.s32 	%rd44, %r52, 2;
+	add.s64 	%rd43, %rd17, %rd44;
+	.loc	1 36 50
+	mov.u32 %r43, 0x0;
+	mov.u32 %r44, 0x0;
+	@%p1 ld.global.L1::evict_last.v2.b32 { %r43, %r44 }, [ %rd43 + 0 ];
+	@!%p1 mov.u32 %r43, %r143;
+	@!%p1 mov.u32 %r44, %r143;
+	mov.b32 	%r155, 883;
+	mov.u64 	%rd67, 1;
+	.loc	1 40 55
+	@%p16 bra 	$L__BB0_3;
+	mov.u64 	%rd45, assertMessage_0;
+	cvta.global.u64 	%rd46, %rd45;
+	mov.u64 	%rd47, assertFile_0;
+	cvta.global.u64 	%rd48, %rd47;
+	mov.u64 	%rd49, assertFunc_0;
+	cvta.global.u64 	%rd50, %rd49;
+	{ // callseq 2, 0
+	.reg .b32 temp_param_reg;
+	.param .b64 param0;
+	st.param.b64 	[param0+0], %rd46;
+	.param .b64 param1;
+	st.param.b64 	[param1+0], %rd48;
+	.param .b32 param2;
+	st.param.b32 	[param2+0], %r155;
+	.param .b64 param3;
+	st.param.b64 	[param3+0], %rd50;
+	.param .b64 param4;
+	st.param.b64 	[param4+0], %rd67;
+	call.uni 
+	__assertfail, 
+	(
+	param0, 
+	param1, 
+	param2, 
+	param3, 
+	param4
+	);
+	} // callseq 2
+	bra.uni 	$L__BB0_3;
+$L__BB0_4:
+	.loc	1 31 36
+	shr.u32 	%r99, %r1, 3;
+	or.b32  	%r100, %r5, %r99;
+	mad.lo.s32 	%r101, %r100, 12, %r4;
+	shl.b32 	%r102, %r101, 2;
+	mov.u32 	%r103, global_smem;
+	add.s32 	%r104, %r103, %r102;
+	st.shared.f32 	[%r104], %f152;
+	st.shared.f32 	[%r104+768], %f153;
+	st.shared.f32 	[%r104+1536], %f154;
+	st.shared.f32 	[%r104+2304], %f155;
+	bar.sync 	0;
+	mad.lo.s32 	%r105, %r2, 12, %r3;
+	shl.b32 	%r106, %r105, 2;
+	add.s32 	%r107, %r103, %r106;
+	ld.shared.v4.f32 	{%f80, %f81, %f82, %f83}, [%r107];
+$L__tmp3:
+	.loc	2 108 21
+	sub.f32 	%f84, %f161, %f160;
+	.loc	2 109 28
+	add.f32 	%f85, %f80, %f81;
+	.loc	2 110 39
+	setp.eq.f32 	%p23, %f85, 0f00000000;
+	.loc	2 110 60
+	mov.b32 	%r75, %f81;
+	mov.b32 	%r76, %f85;
+	div.full.f32 %r74, %r75, %r76;
+	mov.b32 	%f86, %r74;
+	.loc	2 110 49
+	selp.f32 	%f87, 0f00000000, %f86, %p23;
+	.loc	2 112 17
+	fma.rn.f32 	%f88, %f84, %f87, %f160;
+	.loc	2 113 15
+	add.f32 	%f89, %f156, %f157;
+	.loc	2 113 30
+	mul.f32 	%f90, %f84, %f84;
+	.loc	2 113 38
+	mul.f32 	%f91, %f90, %f80;
+	.loc	2 113 22
+	fma.rn.f32 	%f92, %f91, %f87, %f89;
+	.loc	2 108 21
+	sub.f32 	%f93, %f162, %f88;
+	.loc	2 109 28
+	add.f32 	%f94, %f82, %f85;
+	.loc	2 110 39
+	setp.eq.f32 	%p24, %f94, 0f00000000;
+	.loc	2 110 60
+	mov.b32 	%r79, %f94;
+	mov.b32 	%r78, %f82;
+	div.full.f32 %r77, %r78, %r79;
+	mov.b32 	%f95, %r77;
+	.loc	2 110 49
+	selp.f32 	%f96, 0f00000000, %f95, %p24;
+	.loc	2 112 17
+	fma.rn.f32 	%f97, %f96, %f93, %f88;
+	.loc	2 113 15
+	add.f32 	%f98, %f158, %f92;
+	.loc	2 113 30
+	mul.f32 	%f99, %f93, %f93;
+	.loc	2 113 38
+	mul.f32 	%f100, %f85, %f99;
+	.loc	2 113 22
+	fma.rn.f32 	%f101, %f96, %f100, %f98;
+	.loc	2 108 21
+	sub.f32 	%f102, %f163, %f97;
+	.loc	2 109 28
+	add.f32 	%f103, %f83, %f94;
+	.loc	2 110 39
+	setp.eq.f32 	%p25, %f103, 0f00000000;
+	.loc	2 110 60
+	mov.b32 	%r82, %f103;
+	mov.b32 	%r81, %f83;
+	div.full.f32 %r80, %r81, %r82;
+	mov.b32 	%f104, %r80;
+	.loc	2 110 49
+	selp.f32 	%f105, 0f00000000, %f104, %p25;
+	.loc	2 112 17
+	fma.rn.f32 	%f106, %f105, %f102, %f97;
+	.loc	2 113 15
+	add.f32 	%f107, %f159, %f101;
+	.loc	2 113 30
+	mul.f32 	%f108, %f102, %f102;
+	.loc	2 113 38
+	mul.f32 	%f109, %f94, %f108;
+	.loc	2 113 22
+	fma.rn.f32 	%f110, %f105, %f109, %f107;
+$L__tmp4:
+	.loc	2 120 46
+	mov.b32 	%r108, %f106;
+	shfl.sync.bfly.b32	%r109, %r108, 1, 31, -1;
+	mov.b32 	%f111, %r109;
+	mov.b32 	%r110, %f110;
+	shfl.sync.bfly.b32	%r111, %r110, 1, 31, -1;
+	mov.b32 	%f112, %r111;
+	shfl.sync.bfly.b32	%r84, %r82, 1, 31, -1;
+	mov.b32 	%f113, %r84;
+$L__tmp5:
+	.loc	2 108 21
+	sub.f32 	%f114, %f111, %f106;
+	.loc	2 109 28
+	add.f32 	%f115, %f103, %f113;
+	.loc	2 110 39
+	setp.eq.f32 	%p26, %f115, 0f00000000;
+	.loc	2 110 60
+	mov.b32 	%r85, %f115;
+	div.full.f32 %r83, %r84, %r85;
+	mov.b32 	%f116, %r83;
+	.loc	2 110 49
+	selp.f32 	%f117, 0f00000000, %f116, %p26;
+	.loc	2 112 17
+	fma.rn.f32 	%f41, %f117, %f114, %f106;
+	.loc	2 113 15
+	add.f32 	%f118, %f110, %f112;
+	.loc	2 113 30
+	mul.f32 	%f119, %f114, %f114;
+	.loc	2 113 38
+	mul.f32 	%f120, %f103, %f119;
+	.loc	2 113 22
+	fma.rn.f32 	%f121, %f117, %f120, %f118;
+$L__tmp6:
+	.loc	1 75 24
+	mov.b32 	%r87, %f121;
+	mov.b32 	%r88, 1132462080;
+	div.full.f32 %r86, %r87, %r88;
+	mov.b32 	%f122, %r86;
+	.loc	1 77 24
+	add.f32 	%f42, %f122, 0f3727C5AC;
+	.loc	1 58 36
+	add.s64 	%rd71, %rd18, %rd2;
+	mov.b32 	%r157, -8;
+	rsqrt.approx.ftz.f32 	%f139, %f42;
+	bra.uni 	$L__BB0_5;
+$L__BB0_7:
+	.loc	1 0 0
+	mov.b32 	%f43, %r112;
+	mov.b32 	%f44, %r113;
+	mov.b32 	%f45, %r114;
+	mov.b32 	%f46, %r115;
+	cvt.s64.s32 	%rd13, %r137;
+	mov.b32 	%f47, %r124;
+	mov.b32 	%f48, %r125;
+	mov.b32 	%f49, %r126;
+	mov.b32 	%f50, %r127;
+	mov.b32 	%f51, %r128;
+	mov.b32 	%f52, %r129;
+	mov.b32 	%f53, %r130;
+	mov.b32 	%f54, %r131;
+	.loc	1 69 54
+	mov.u32 %r139, 0x0;
+	mov.u32 %r140, 0x0;
+	mov.u32 %r141, 0x0;
+	mov.u32 %r142, 0x0;
+	@%p1 ld.global.L1::evict_first.v4.b32 { %r139, %r140, %r141, %r142 }, [ %rd72 + 0 ];
+	@!%p1 mov.u32 %r139, %r143;
+	@!%p1 mov.u32 %r140, %r143;
+	@!%p1 mov.u32 %r141, %r143;
+	@!%p1 mov.u32 %r142, %r143;
+	mov.b32 	%f123, %r139;
+	mov.b32 	%f124, %r140;
+	mov.b32 	%f125, %r141;
+	mov.b32 	%f126, %r142;
+	.loc	1 70 24
+	add.f32 	%f127, %f43, %f123;
+	add.f32 	%f128, %f44, %f124;
+	add.f32 	%f129, %f45, %f125;
+	add.f32 	%f130, %f46, %f126;
+	.loc	1 72 24
+	add.f32 	%f131, %f47, %f127;
+	add.f32 	%f132, %f48, %f128;
+	add.f32 	%f133, %f49, %f129;
+	add.f32 	%f134, %f50, %f130;
+	.loc	1 73 24
+	sub.f32 	%f135, %f131, %f41;
+	sub.f32 	%f136, %f132, %f41;
+	sub.f32 	%f137, %f133, %f41;
+	sub.f32 	%f138, %f134, %f41;
+	.loc	1 79 24
+	mul.f32 	%f140, %f135, %f139;
+	mul.f32 	%f141, %f136, %f139;
+	mul.f32 	%f142, %f137, %f139;
+	mul.f32 	%f143, %f138, %f139;
+	.loc	1 80 24
+	mul.f32 	%f144, %f140, %f51;
+	mul.f32 	%f145, %f141, %f52;
+	mul.f32 	%f146, %f142, %f53;
+	mul.f32 	%f147, %f143, %f54;
+	.loc	1 82 29
+	shl.b64 	%rd66, %rd13, 1;
+	add.s64 	%rd65, %rd19, %rd66;
+	.loc	1 82 52
+	mov.b32 	%r147, %f144;
+	cvt.rn.bf16.f32 %rs9, %r147;
+	mov.b32 	%r148, %f145;
+	cvt.rn.bf16.f32 %rs10, %r148;
+	mov.b32 	%r149, %f146;
+	cvt.rn.bf16.f32 %rs11, %r149;
+	mov.b32 	%r150, %f147;
+	cvt.rn.bf16.f32 %rs12, %r150;
+	mov.b32 	%r153, {%rs9, %rs10};
+	mov.b32 	%r154, {%rs11, %rs12};
+	@%p1 st.global.v2.b32 [ %rd65 + 0 ], { %r153, %r154 };
+	.loc	1 58 36
+	add.s32 	%r157, %r157, 8;
+	add.s64 	%rd72, %rd72, 32;
+	add.s64 	%rd71, %rd71, 32;
+	add.s64 	%rd70, %rd70, 32;
+	setp.lt.u32 	%p47, %r157, 248;
+	@%p47 bra 	$L__BB0_5;
+	bra.uni 	$L__BB0_8;
+$L__BB0_5:
+	.loc	1 62 51
+	mov.u32 %r112, 0x0;
+	mov.u32 %r113, 0x0;
+	mov.u32 %r114, 0x0;
+	mov.u32 %r115, 0x0;
+	@%p1 ld.global.L1::evict_last.v4.b32 { %r112, %r113, %r114, %r115 }, [ %rd70 + 0 ];
+	@!%p1 mov.u32 %r112, %r143;
+	@!%p1 mov.u32 %r113, %r143;
+	@!%p1 mov.u32 %r114, %r143;
+	@!%p1 mov.u32 %r115, %r143;
+	.loc	1 63 35
+	add.s32 	%r136, %r6, %r157;
+	add.s32 	%r137, %r136, 8;
+	mul.wide.s32 	%rd56, %r137, 2;
+	add.s64 	%rd54, %rd17, %rd56;
+	.loc	1 63 51
+	mov.u32 %r120, 0x0;
+	mov.u32 %r121, 0x0;
+	@%p1 ld.global.L1::evict_first.v2.b32 { %r120, %r121 }, [ %rd54 + 0 ];
+	@!%p1 mov.u32 %r120, %r143;
+	@!%p1 mov.u32 %r121, %r143;
+	cvt.u16.u32 	%rs5, %r120;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r120; }
+	cvt.u16.u32 	%rs7, %r121;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r121; }
+	.loc	1 63 103
+	cvt.f32.bf16 %r124, %rs5;
+	cvt.f32.bf16 %r125, %rs6;
+	cvt.f32.bf16 %r126, %rs7;
+	cvt.f32.bf16 %r127, %rs8;
+	.loc	1 64 40
+	mov.u32 %r128, 0x0;
+	mov.u32 %r129, 0x0;
+	mov.u32 %r130, 0x0;
+	mov.u32 %r131, 0x0;
+	@%p1 ld.global.L1::evict_last.v4.b32 { %r128, %r129, %r130, %r131 }, [ %rd71 + 0 ];
+	@!%p1 mov.u32 %r128, %r143;
+	@!%p1 mov.u32 %r129, %r143;
+	@!%p1 mov.u32 %r130, %r143;
+	@!%p1 mov.u32 %r131, %r143;
+	.loc	1 68 57
+	@%p16 bra 	$L__BB0_7;
+	mov.u64 	%rd57, assertMessage_1;
+	cvta.global.u64 	%rd58, %rd57;
+	mov.u64 	%rd59, assertFile_1;
+	cvta.global.u64 	%rd60, %rd59;
+	mov.u64 	%rd61, assertFunc_1;
+	cvta.global.u64 	%rd62, %rd61;
+	{ // callseq 3, 0
+	.reg .b32 temp_param_reg;
+	.param .b64 param0;
+	st.param.b64 	[param0+0], %rd58;
+	.param .b64 param1;
+	st.param.b64 	[param1+0], %rd60;
+	.param .b32 param2;
+	st.param.b32 	[param2+0], %r155;
+	.param .b64 param3;
+	st.param.b64 	[param3+0], %rd62;
+	.param .b64 param4;
+	st.param.b64 	[param4+0], %rd67;
+	call.uni 
+	__assertfail, 
+	(
+	param0, 
+	param1, 
+	param2, 
+	param3, 
+	param4
+	);
+	} // callseq 3
+	bra.uni 	$L__BB0_7;
+$L__BB0_8:
+	.loc	1 58 4
+	ret;
+$L__tmp7:
+$L__func_end0:
+
+}
+	// .globl	__nv_rsqrtf
+.visible .func  (.param .b32 func_retval0) __nv_rsqrtf(
+	.param .b32 __nv_rsqrtf_param_0
+)
+{
+	.reg .f32 	%f<3>;
+$L__func_begin1:
+
+	ld.param.f32 	%f1, [__nv_rsqrtf_param_0];
+	rsqrt.approx.ftz.f32 	%f2, %f1;
+	st.param.f32 	[func_retval0+0], %f2;
+	ret;
+$L__func_end1:
+
+}
+	.file	1 "/tmp/torchinductor_root/pn/cpn3lawg65lpi63gv6c6pn4oikhg6qva2h2qjdpxe6qj4lvttwez.py"
+	.file	2 "/usr/local/lib/python3.10/dist-packages/torch/_inductor/triton_helpers.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 302
+.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 112
+.b8 110
+.b8 51
+.b8 108
+.b8 97
+.b8 119
+.b8 103
+.b8 54
+.b8 53
+.b8 108
+.b8 112
+.b8 105
+.b8 54
+.b8 51
+.b8 103
+.b8 118
+.b8 54
+.b8 99
+.b8 54
+.b8 112
+.b8 110
+.b8 52
+.b8 111
+.b8 105
+.b8 107
+.b8 104
+.b8 103
+.b8 54
+.b8 113
+.b8 118
+.b8 97
+.b8 50
+.b8 104
+.b8 50
+.b8 113
+.b8 106
+.b8 100
+.b8 112
+.b8 120
+.b8 101
+.b8 54
+.b8 113
+.b8 106
+.b8 52
+.b8 108
+.b8 118
+.b8 116
+.b8 116
+.b8 119
+.b8 101
+.b8 122
+.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 112
+.b8 110
+.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 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 54
+.b8 100
+.b8 101
+.b8 55
+.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 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 54
+.b8 100
+.b8 101
+.b8 55
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 1
+.b8 3
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b32 125
+.b8 4
+.b32 125
+.b64 $L__tmp1
+.b64 $L__tmp2
+.b8 2
+.b8 47
+.b8 41
+.b8 5
+.b32 125
+.b64 $L__tmp3
+.b64 $L__tmp6
+.b8 2
+.b8 53
+.b8 44
+.b8 4
+.b32 125
+.b64 $L__tmp3
+.b64 $L__tmp6
+.b8 2
+.b8 120
+.b8 46
+.b8 0
+.b8 4
+.b32 125
+.b64 $L__tmp4
+.b64 $L__tmp5
+.b8 2
+.b8 53
+.b8 44
+.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 306
+.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 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 54
+.b8 100
+.b8 101
+.b8 55
+.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 306
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}

.triton/dump/0471aff594c8c8b8715b81c529738739/triton_.ttir

@@ -0,0 +1,153 @@
+module {
+  tt.func public @triton__0d1d2d3d4d5d6de7de(%arg0: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg4: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg5: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg6: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}, %arg7: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<0.000000e+00> : tensor<64x8xbf16>
+    %cst_0 = arith.constant 0.000000e+00 : f32
+    %cst_1 = arith.constant dense<1.000000e+00> : tensor<64x8xf32>
+    %c256_i32 = arith.constant 256 : i32
+    %c8_i32 = arith.constant 8 : i32
+    %c0_i32 = arith.constant 0 : i32
+    %cst_2 = arith.constant dense<256> : tensor<64x1xi64>
+    %cst_3 = arith.constant dense<0> : tensor<64x1xi64>
+    %cst_4 = arith.constant dense<50257> : tensor<64x1xi64>
+    %cst_5 = arith.constant dense<9.99999974E-6> : tensor<64x1xf32>
+    %cst_6 = arith.constant dense<2.560000e+02> : tensor<64x1xf32>
+    %cst_7 = arith.constant dense<0.000000e+00> : tensor<1x8xf32>
+    %cst_8 = arith.constant dense<0.000000e+00> : tensor<64x8xf32>
+    %cst_9 = arith.constant dense<256> : tensor<64x1xi32>
+    %cst_10 = arith.constant dense<256> : tensor<1x8xi32>
+    %cst_11 = arith.constant dense<512> : tensor<64x1xi32>
+    %c64_i32 = arith.constant 64 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c64_i32 : i32
+    %2 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32>
+    %3 = tt.expand_dims %2 {axis = 1 : i32} : (tensor<64xi32>) -> tensor<64x1xi32>
+    %4 = tt.splat %1 : (i32) -> tensor<64x1xi32>
+    %5 = arith.addi %4, %3 : tensor<64x1xi32>
+    %6 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32>
+    %7 = tt.expand_dims %6 {axis = 0 : i32} : (tensor<8xi32>) -> tensor<1x8xi32>
+    %8 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<64x1x!tt.ptr<i64, 1>>
+    %9 = tt.addptr %8, %5 : tensor<64x1x!tt.ptr<i64, 1>>, tensor<64x1xi32>
+    %10 = tt.load %9 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x1xi64>
+    %11 = arith.remsi %5, %cst_11 : tensor<64x1xi32>
+    %12 = arith.muli %11, %cst_9 : tensor<64x1xi32>
+    %13 = tt.broadcast %12 : (tensor<64x1xi32>) -> tensor<64x8xi32>
+    %14 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>>
+    %15 = arith.muli %5, %cst_9 : tensor<64x1xi32>
+    %16 = tt.broadcast %15 : (tensor<64x1xi32>) -> tensor<64x8xi32>
+    %17 = tt.splat %arg3 : (!tt.ptr<bf16, 1>) -> tensor<64x8x!tt.ptr<bf16, 1>>
+    %18 = arith.addi %10, %cst_4 : tensor<64x1xi64>
+    %19 = arith.cmpi slt, %10, %cst_3 : tensor<64x1xi64>
+    %20 = arith.select %19, %18, %10 : tensor<64x1xi1>, tensor<64x1xi64>
+    %21 = arith.cmpi sge, %20, %cst_3 : tensor<64x1xi64>
+    %22 = arith.cmpi slt, %20, %cst_4 : tensor<64x1xi64>
+    %23 = arith.andi %21, %22 : tensor<64x1xi1>
+    %24 = arith.muli %20, %cst_2 : tensor<64x1xi64>
+    %25 = tt.broadcast %24 : (tensor<64x1xi64>) -> tensor<64x8xi64>
+    %26 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>>
+    %27:3 = scf.for %arg8 = %c0_i32 to %c256_i32 step %c8_i32 iter_args(%arg9 = %cst_8, %arg10 = %cst_8, %arg11 = %cst_8) -> (tensor<64x8xf32>, tensor<64x8xf32>, tensor<64x8xf32>)  : i32 {
+      %51 = tt.splat %arg8 : (i32) -> tensor<1x8xi32>
+      %52 = arith.addi %51, %7 : tensor<1x8xi32>
+      %53 = arith.cmpi slt, %52, %cst_10 : tensor<1x8xi32>
+      %54 = tt.broadcast %52 : (tensor<1x8xi32>) -> tensor<64x8xi32>
+      %55 = arith.addi %54, %13 : tensor<64x8xi32>
+      %56 = tt.addptr %14, %55 : tensor<64x8x!tt.ptr<f32, 1>>, tensor<64x8xi32>
+      %57 = tt.broadcast %53 : (tensor<1x8xi1>) -> tensor<64x8xi1>
+      %58 = tt.load %56, %57, %cst_8 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xf32>
+      %59 = arith.addi %54, %16 : tensor<64x8xi32>
+      %60 = tt.addptr %17, %59 : tensor<64x8x!tt.ptr<bf16, 1>>, tensor<64x8xi32>
+      %61 = tt.load %60, %57, %cst {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xbf16>
+      %62 = arith.extf %61 : tensor<64x8xbf16> to tensor<64x8xf32>
+      tt.assert %23, "index out of bounds: 0 <= tmp3 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<64x1xi1>
+      %63 = arith.extsi %52 : tensor<1x8xi32> to tensor<1x8xi64>
+      %64 = tt.broadcast %63 : (tensor<1x8xi64>) -> tensor<64x8xi64>
+      %65 = arith.addi %64, %25 : tensor<64x8xi64>
+      %66 = tt.addptr %26, %65 : tensor<64x8x!tt.ptr<f32, 1>>, tensor<64x8xi64>
+      %67 = tt.load %66, %57, %cst_8 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xf32>
+      %68 = arith.addf %67, %58 : tensor<64x8xf32>
+      %69 = arith.addf %68, %62 : tensor<64x8xf32>
+      %70 = arith.subf %69, %arg9 : tensor<64x8xf32>
+      %71 = arith.addf %arg11, %cst_1 : tensor<64x8xf32>
+      %72 = arith.divf %70, %71 : tensor<64x8xf32>
+      %73 = arith.addf %arg9, %72 : tensor<64x8xf32>
+      %74 = arith.subf %69, %73 : tensor<64x8xf32>
+      %75 = arith.mulf %70, %74 : tensor<64x8xf32>
+      %76 = arith.addf %arg10, %75 : tensor<64x8xf32>
+      %77 = arith.select %57, %73, %arg9 : tensor<64x8xi1>, tensor<64x8xf32>
+      %78 = arith.select %57, %76, %arg10 : tensor<64x8xi1>, tensor<64x8xf32>
+      %79 = arith.select %57, %71, %arg11 : tensor<64x8xi1>, tensor<64x8xf32>
+      scf.yield %77, %78, %79 : tensor<64x8xf32>, tensor<64x8xf32>, tensor<64x8xf32>
+    }
+    %28:3 = "tt.reduce"(%27#0, %27#1, %27#2) <{axis = 1 : i32}> ({
+    ^bb0(%arg8: f32, %arg9: f32, %arg10: f32, %arg11: f32, %arg12: f32, %arg13: f32):
+      %51 = arith.subf %arg11, %arg8 : f32
+      %52 = arith.addf %arg10, %arg13 : f32
+      %53 = arith.cmpf oeq, %52, %cst_0 : f32
+      %54 = arith.divf %arg13, %52 : f32
+      %55 = arith.select %53, %cst_0, %54 : f32
+      %56 = arith.mulf %51, %55 : f32
+      %57 = arith.addf %arg8, %56 : f32
+      %58 = arith.addf %arg9, %arg12 : f32
+      %59 = arith.mulf %51, %51 : f32
+      %60 = arith.mulf %59, %arg10 : f32
+      %61 = arith.mulf %60, %55 : f32
+      %62 = arith.addf %58, %61 : f32
+      tt.reduce.return %57, %62, %52 : f32, f32, f32
+    }) : (tensor<64x8xf32>, tensor<64x8xf32>, tensor<64x8xf32>) -> (tensor<64xf32>, tensor<64xf32>, tensor<64xf32>)
+    %29 = tt.expand_dims %28#0 {axis = 1 : i32} : (tensor<64xf32>) -> tensor<64x1xf32>
+    %30 = tt.expand_dims %28#1 {axis = 1 : i32} : (tensor<64xf32>) -> tensor<64x1xf32>
+    %31 = arith.muli %11, %cst_9 : tensor<64x1xi32>
+    %32 = tt.broadcast %31 : (tensor<64x1xi32>) -> tensor<64x8xi32>
+    %33 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>>
+    %34 = arith.muli %5, %cst_9 : tensor<64x1xi32>
+    %35 = tt.broadcast %34 : (tensor<64x1xi32>) -> tensor<64x8xi32>
+    %36 = tt.splat %arg3 : (!tt.ptr<bf16, 1>) -> tensor<64x8x!tt.ptr<bf16, 1>>
+    %37 = tt.splat %arg4 : (!tt.ptr<f32, 1>) -> tensor<1x8x!tt.ptr<f32, 1>>
+    %38 = arith.addi %10, %cst_4 : tensor<64x1xi64>
+    %39 = arith.cmpi slt, %10, %cst_3 : tensor<64x1xi64>
+    %40 = arith.select %39, %38, %10 : tensor<64x1xi1>, tensor<64x1xi64>
+    %41 = arith.cmpi sge, %40, %cst_3 : tensor<64x1xi64>
+    %42 = arith.cmpi slt, %40, %cst_4 : tensor<64x1xi64>
+    %43 = arith.andi %41, %42 : tensor<64x1xi1>
+    %44 = arith.muli %40, %cst_2 : tensor<64x1xi64>
+    %45 = tt.broadcast %44 : (tensor<64x1xi64>) -> tensor<64x8xi64>
+    %46 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>>
+    %47 = tt.broadcast %29 : (tensor<64x1xf32>) -> tensor<64x8xf32>
+    %48 = arith.divf %30, %cst_6 : tensor<64x1xf32>
+    %49 = arith.addf %48, %cst_5 : tensor<64x1xf32>
+    %50 = tt.splat %arg5 : (!tt.ptr<bf16, 1>) -> tensor<64x8x!tt.ptr<bf16, 1>>
+    scf.for %arg8 = %c0_i32 to %c256_i32 step %c8_i32  : i32 {
+      %51 = tt.splat %arg8 : (i32) -> tensor<1x8xi32>
+      %52 = arith.addi %51, %7 : tensor<1x8xi32>
+      %53 = arith.cmpi slt, %52, %cst_10 : tensor<1x8xi32>
+      %54 = tt.broadcast %52 : (tensor<1x8xi32>) -> tensor<64x8xi32>
+      %55 = arith.addi %54, %32 : tensor<64x8xi32>
+      %56 = tt.addptr %33, %55 : tensor<64x8x!tt.ptr<f32, 1>>, tensor<64x8xi32>
+      %57 = tt.broadcast %53 : (tensor<1x8xi1>) -> tensor<64x8xi1>
+      %58 = tt.load %56, %57, %cst_8 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xf32>
+      %59 = arith.addi %54, %35 : tensor<64x8xi32>
+      %60 = tt.addptr %36, %59 : tensor<64x8x!tt.ptr<bf16, 1>>, tensor<64x8xi32>
+      %61 = tt.load %60, %57, %cst {cache = 1 : i32, evict = 2 : i32, isVolatile = false} : tensor<64x8xbf16>
+      %62 = arith.extf %61 : tensor<64x8xbf16> to tensor<64x8xf32>
+      %63 = tt.addptr %37, %52 : tensor<1x8x!tt.ptr<f32, 1>>, tensor<1x8xi32>
+      %64 = tt.load %63, %53, %cst_7 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<1x8xf32>
+      tt.assert %43, "index out of bounds: 0 <= tmp16 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<64x1xi1>
+      %65 = arith.extsi %52 : tensor<1x8xi32> to tensor<1x8xi64>
+      %66 = tt.broadcast %65 : (tensor<1x8xi64>) -> tensor<64x8xi64>
+      %67 = arith.addi %66, %45 : tensor<64x8xi64>
+      %68 = tt.addptr %46, %67 : tensor<64x8x!tt.ptr<f32, 1>>, tensor<64x8xi64>
+      %69 = tt.load %68, %57, %cst_8 {cache = 1 : i32, evict = 2 : i32, isVolatile = false} : tensor<64x8xf32>
+      %70 = arith.addf %69, %58 : tensor<64x8xf32>
+      %71 = arith.addf %70, %62 : tensor<64x8xf32>
+      %72 = arith.subf %71, %47 : tensor<64x8xf32>
+      %73 = tt.extern_elementwise %49 {libname = "libdevice", libpath = "/usr/local/lib/python3.10/dist-packages/triton/language/../third_party/cuda/lib/libdevice.10.bc", pure = true, symbol = "__nv_rsqrtf"} : (tensor<64x1xf32>) -> tensor<64x1xf32>
+      %74 = tt.broadcast %73 : (tensor<64x1xf32>) -> tensor<64x8xf32>
+      %75 = arith.mulf %72, %74 : tensor<64x8xf32>
+      %76 = tt.broadcast %64 : (tensor<1x8xf32>) -> tensor<64x8xf32>
+      %77 = arith.mulf %75, %76 : tensor<64x8xf32>
+      %78 = tt.addptr %50, %59 : tensor<64x8x!tt.ptr<bf16, 1>>, tensor<64x8xi32>
+      %79 = arith.truncf %77 : tensor<64x8xf32> to tensor<64x8xbf16>
+      tt.store %78, %79, %57 {cache = 1 : i32, evict = 1 : i32} : tensor<64x8xbf16>
+    }
+    tt.return
+  }
+}

.triton/dump/9f68cc707cb8f8bff3232abf59cbd9ec/triton_.ttgir

@@ -0,0 +1,154 @@
+#blocked = #triton_gpu.blocked<{sizePerThread = [1, 4], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [1, 0], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [1, 0]}>
+#blocked1 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 1], order = [1, 0], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [1, 0]}>
+#blocked2 = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [4, 8], warpsPerCTA = [4, 1], order = [1, 0], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [1, 0]}>
+module attributes {"triton_gpu.compute-capability" = 89 : i32, "triton_gpu.num-ctas" = 1 : i32, "triton_gpu.num-warps" = 4 : i32, "triton_gpu.threads-per-warp" = 32 : i32} {
+  tt.func public @triton__0d1d2d3d4d5de6de(%arg0: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg4: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}, %arg6: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<512> : tensor<64x1xi32, #blocked>
+    %cst_0 = arith.constant dense<256> : tensor<1x8xi32, #blocked>
+    %cst_1 = arith.constant dense<256> : tensor<64x1xi32, #blocked>
+    %cst_2 = arith.constant dense<0.000000e+00> : tensor<64x8xf32, #blocked>
+    %cst_3 = arith.constant dense<0.000000e+00> : tensor<1x8xf32, #blocked>
+    %cst_4 = arith.constant dense<1.000000e+00> : tensor<64x8xf32, #blocked>
+    %cst_5 = arith.constant dense<256> : tensor<64x1xi64, #blocked>
+    %cst_6 = arith.constant dense<0> : tensor<64x1xi64, #blocked>
+    %cst_7 = arith.constant dense<50257> : tensor<64x1xi64, #blocked>
+    %cst_8 = arith.constant dense<50257> : tensor<64x1xi64, #blocked1>
+    %cst_9 = arith.constant dense<0> : tensor<64x1xi64, #blocked1>
+    %c0_i32 = arith.constant 0 : i32
+    %c8_i32 = arith.constant 8 : i32
+    %c256_i32 = arith.constant 256 : i32
+    %cst_10 = arith.constant dense<1.000000e+00> : tensor<64x8xf32, #blocked2>
+    %cst_11 = arith.constant 0.000000e+00 : f32
+    %cst_12 = arith.constant dense<0.000000e+00> : tensor<64x8xf32, #blocked2>
+    %cst_13 = arith.constant dense<256> : tensor<1x8xi32, #blocked2>
+    %cst_14 = arith.constant dense<9.99999974E-6> : tensor<64x1xf32, #blocked>
+    %cst_15 = arith.constant dense<2.560000e+02> : tensor<64x1xf32, #blocked>
+    %c64_i32 = arith.constant 64 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c64_i32 : i32
+    %2 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>
+    %3 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>
+    %4 = tt.expand_dims %2 {axis = 1 : i32} : (tensor<64xi32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>) -> tensor<64x1xi32, #blocked>
+    %5 = tt.expand_dims %3 {axis = 1 : i32} : (tensor<64xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>) -> tensor<64x1xi32, #blocked1>
+    %6 = tt.splat %1 : (i32) -> tensor<64x1xi32, #blocked>
+    %7 = tt.splat %1 : (i32) -> tensor<64x1xi32, #blocked1>
+    %8 = arith.addi %6, %4 : tensor<64x1xi32, #blocked>
+    %9 = arith.addi %7, %5 : tensor<64x1xi32, #blocked1>
+    %10 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #triton_gpu.slice<{dim = 0, parent = #blocked}>>
+    %11 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #triton_gpu.slice<{dim = 0, parent = #blocked2}>>
+    %12 = tt.expand_dims %10 {axis = 0 : i32} : (tensor<8xi32, #triton_gpu.slice<{dim = 0, parent = #blocked}>>) -> tensor<1x8xi32, #blocked>
+    %13 = tt.expand_dims %11 {axis = 0 : i32} : (tensor<8xi32, #triton_gpu.slice<{dim = 0, parent = #blocked2}>>) -> tensor<1x8xi32, #blocked2>
+    %14 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<64x1x!tt.ptr<i64, 1>, #blocked>
+    %15 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<64x1x!tt.ptr<i64, 1>, #blocked1>
+    %16 = tt.addptr %14, %8 : tensor<64x1x!tt.ptr<i64, 1>, #blocked>, tensor<64x1xi32, #blocked>
+    %17 = tt.addptr %15, %9 : tensor<64x1x!tt.ptr<i64, 1>, #blocked1>, tensor<64x1xi32, #blocked1>
+    %18 = tt.load %16 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x1xi64, #blocked>
+    %19 = tt.load %17 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x1xi64, #blocked1>
+    %20 = arith.remsi %8, %cst : tensor<64x1xi32, #blocked>
+    %21 = arith.muli %20, %cst_1 : tensor<64x1xi32, #blocked>
+    %22 = tt.broadcast %21 : (tensor<64x1xi32, #blocked>) -> tensor<64x8xi32, #blocked>
+    %23 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>, #blocked>
+    %24 = arith.addi %18, %cst_7 : tensor<64x1xi64, #blocked>
+    %25 = arith.addi %19, %cst_8 : tensor<64x1xi64, #blocked1>
+    %26 = arith.cmpi slt, %18, %cst_6 : tensor<64x1xi64, #blocked>
+    %27 = arith.cmpi slt, %19, %cst_9 : tensor<64x1xi64, #blocked1>
+    %28 = arith.select %26, %24, %18 : tensor<64x1xi1, #blocked>, tensor<64x1xi64, #blocked>
+    %29 = arith.select %27, %25, %19 : tensor<64x1xi1, #blocked1>, tensor<64x1xi64, #blocked1>
+    %30 = arith.cmpi sge, %29, %cst_9 : tensor<64x1xi64, #blocked1>
+    %31 = arith.cmpi slt, %29, %cst_8 : tensor<64x1xi64, #blocked1>
+    %32 = arith.andi %30, %31 : tensor<64x1xi1, #blocked1>
+    %33 = arith.muli %28, %cst_5 : tensor<64x1xi64, #blocked>
+    %34 = tt.broadcast %33 : (tensor<64x1xi64, #blocked>) -> tensor<64x8xi64, #blocked>
+    %35 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>, #blocked>
+    %36:4 = scf.for %arg7 = %c0_i32 to %c256_i32 step %c8_i32 iter_args(%arg8 = %cst_2, %arg9 = %cst_2, %arg10 = %cst_12, %arg11 = %cst_2) -> (tensor<64x8xf32, #blocked>, tensor<64x8xf32, #blocked>, tensor<64x8xf32, #blocked2>, tensor<64x8xf32, #blocked>)  : i32 {
+      %48 = tt.splat %arg7 : (i32) -> tensor<1x8xi32, #blocked>
+      %49 = tt.splat %arg7 : (i32) -> tensor<1x8xi32, #blocked2>
+      %50 = arith.addi %48, %12 : tensor<1x8xi32, #blocked>
+      %51 = arith.addi %49, %13 : tensor<1x8xi32, #blocked2>
+      %52 = arith.cmpi slt, %50, %cst_0 : tensor<1x8xi32, #blocked>
+      %53 = arith.cmpi slt, %51, %cst_13 : tensor<1x8xi32, #blocked2>
+      %54 = tt.broadcast %50 : (tensor<1x8xi32, #blocked>) -> tensor<64x8xi32, #blocked>
+      %55 = arith.addi %54, %22 : tensor<64x8xi32, #blocked>
+      %56 = tt.addptr %23, %55 : tensor<64x8x!tt.ptr<f32, 1>, #blocked>, tensor<64x8xi32, #blocked>
+      %57 = tt.broadcast %52 : (tensor<1x8xi1, #blocked>) -> tensor<64x8xi1, #blocked>
+      %58 = tt.broadcast %53 : (tensor<1x8xi1, #blocked2>) -> tensor<64x8xi1, #blocked2>
+      %59 = tt.load %56, %57, %cst_2 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xf32, #blocked>
+      tt.assert %32, "index out of bounds: 0 <= tmp3 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<64x1xi1, #blocked1>
+      %60 = arith.extsi %50 : tensor<1x8xi32, #blocked> to tensor<1x8xi64, #blocked>
+      %61 = tt.broadcast %60 : (tensor<1x8xi64, #blocked>) -> tensor<64x8xi64, #blocked>
+      %62 = arith.addi %61, %34 : tensor<64x8xi64, #blocked>
+      %63 = tt.addptr %35, %62 : tensor<64x8x!tt.ptr<f32, 1>, #blocked>, tensor<64x8xi64, #blocked>
+      %64 = tt.load %63, %57, %cst_2 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xf32, #blocked>
+      %65 = arith.addf %64, %59 : tensor<64x8xf32, #blocked>
+      %66 = arith.subf %65, %arg8 : tensor<64x8xf32, #blocked>
+      %67 = arith.addf %arg11, %cst_4 : tensor<64x8xf32, #blocked>
+      %68 = arith.addf %arg10, %cst_10 : tensor<64x8xf32, #blocked2>
+      %69 = arith.divf %66, %67 : tensor<64x8xf32, #blocked>
+      %70 = arith.addf %arg8, %69 : tensor<64x8xf32, #blocked>
+      %71 = arith.subf %65, %70 : tensor<64x8xf32, #blocked>
+      %72 = arith.mulf %66, %71 : tensor<64x8xf32, #blocked>
+      %73 = arith.addf %arg9, %72 : tensor<64x8xf32, #blocked>
+      %74 = arith.select %57, %70, %arg8 : tensor<64x8xi1, #blocked>, tensor<64x8xf32, #blocked>
+      %75 = arith.select %57, %73, %arg9 : tensor<64x8xi1, #blocked>, tensor<64x8xf32, #blocked>
+      %76 = arith.select %57, %67, %arg11 : tensor<64x8xi1, #blocked>, tensor<64x8xf32, #blocked>
+      %77 = arith.select %58, %68, %arg10 : tensor<64x8xi1, #blocked2>, tensor<64x8xf32, #blocked2>
+      scf.yield %74, %75, %77, %76 : tensor<64x8xf32, #blocked>, tensor<64x8xf32, #blocked>, tensor<64x8xf32, #blocked2>, tensor<64x8xf32, #blocked>
+    }
+    %37 = triton_gpu.convert_layout %36#2 : (tensor<64x8xf32, #blocked2>) -> tensor<64x8xf32, #blocked>
+    %38:3 = "tt.reduce"(%36#0, %36#1, %37) <{axis = 1 : i32}> ({
+    ^bb0(%arg7: f32, %arg8: f32, %arg9: f32, %arg10: f32, %arg11: f32, %arg12: f32):
+      %48 = arith.subf %arg10, %arg7 : f32
+      %49 = arith.addf %arg9, %arg12 : f32
+      %50 = arith.cmpf oeq, %49, %cst_11 : f32
+      %51 = arith.divf %arg12, %49 : f32
+      %52 = arith.select %50, %cst_11, %51 : f32
+      %53 = arith.mulf %48, %52 : f32
+      %54 = arith.addf %arg7, %53 : f32
+      %55 = arith.addf %arg8, %arg11 : f32
+      %56 = arith.mulf %48, %48 : f32
+      %57 = arith.mulf %56, %arg9 : f32
+      %58 = arith.mulf %57, %52 : f32
+      %59 = arith.addf %55, %58 : f32
+      tt.reduce.return %54, %59, %49 : f32, f32, f32
+    }) : (tensor<64x8xf32, #blocked>, tensor<64x8xf32, #blocked>, tensor<64x8xf32, #blocked>) -> (tensor<64xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>, tensor<64xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>, tensor<64xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>)
+    %39 = tt.expand_dims %38#0 {axis = 1 : i32} : (tensor<64xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>) -> tensor<64x1xf32, #blocked>
+    %40 = tt.expand_dims %38#1 {axis = 1 : i32} : (tensor<64xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>) -> tensor<64x1xf32, #blocked>
+    %41 = tt.splat %arg3 : (!tt.ptr<f32, 1>) -> tensor<1x8x!tt.ptr<f32, 1>, #blocked>
+    %42 = tt.broadcast %39 : (tensor<64x1xf32, #blocked>) -> tensor<64x8xf32, #blocked>
+    %43 = arith.divf %40, %cst_15 : tensor<64x1xf32, #blocked>
+    %44 = arith.addf %43, %cst_14 : tensor<64x1xf32, #blocked>
+    %45 = arith.muli %8, %cst_1 : tensor<64x1xi32, #blocked>
+    %46 = tt.broadcast %45 : (tensor<64x1xi32, #blocked>) -> tensor<64x8xi32, #blocked>
+    %47 = tt.splat %arg4 : (!tt.ptr<bf16, 1>) -> tensor<64x8x!tt.ptr<bf16, 1>, #blocked>
+    scf.for %arg7 = %c0_i32 to %c256_i32 step %c8_i32  : i32 {
+      %48 = tt.splat %arg7 : (i32) -> tensor<1x8xi32, #blocked>
+      %49 = arith.addi %48, %12 : tensor<1x8xi32, #blocked>
+      %50 = arith.cmpi slt, %49, %cst_0 : tensor<1x8xi32, #blocked>
+      %51 = tt.broadcast %49 : (tensor<1x8xi32, #blocked>) -> tensor<64x8xi32, #blocked>
+      %52 = arith.addi %51, %22 : tensor<64x8xi32, #blocked>
+      %53 = tt.addptr %23, %52 : tensor<64x8x!tt.ptr<f32, 1>, #blocked>, tensor<64x8xi32, #blocked>
+      %54 = tt.broadcast %50 : (tensor<1x8xi1, #blocked>) -> tensor<64x8xi1, #blocked>
+      %55 = tt.load %53, %54, %cst_2 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<64x8xf32, #blocked>
+      %56 = tt.addptr %41, %49 : tensor<1x8x!tt.ptr<f32, 1>, #blocked>, tensor<1x8xi32, #blocked>
+      %57 = tt.load %56, %50, %cst_3 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<1x8xf32, #blocked>
+      tt.assert %32, "index out of bounds: 0 <= tmp13 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<64x1xi1, #blocked1>
+      %58 = arith.extsi %49 : tensor<1x8xi32, #blocked> to tensor<1x8xi64, #blocked>
+      %59 = tt.broadcast %58 : (tensor<1x8xi64, #blocked>) -> tensor<64x8xi64, #blocked>
+      %60 = arith.addi %59, %34 : tensor<64x8xi64, #blocked>
+      %61 = tt.addptr %35, %60 : tensor<64x8x!tt.ptr<f32, 1>, #blocked>, tensor<64x8xi64, #blocked>
+      %62 = tt.load %61, %54, %cst_2 {cache = 1 : i32, evict = 2 : i32, isVolatile = false} : tensor<64x8xf32, #blocked>
+      %63 = arith.addf %62, %55 : tensor<64x8xf32, #blocked>
+      %64 = arith.subf %63, %42 : tensor<64x8xf32, #blocked>
+      %65 = tt.extern_elementwise %44 {libname = "libdevice", libpath = "/usr/local/lib/python3.10/dist-packages/triton/language/../third_party/cuda/lib/libdevice.10.bc", pure = true, symbol = "__nv_rsqrtf"} : (tensor<64x1xf32, #blocked>) -> tensor<64x1xf32, #blocked>
+      %66 = tt.broadcast %65 : (tensor<64x1xf32, #blocked>) -> tensor<64x8xf32, #blocked>
+      %67 = arith.mulf %64, %66 : tensor<64x8xf32, #blocked>
+      %68 = tt.broadcast %57 : (tensor<1x8xf32, #blocked>) -> tensor<64x8xf32, #blocked>
+      %69 = arith.mulf %67, %68 : tensor<64x8xf32, #blocked>
+      %70 = arith.addi %51, %46 : tensor<64x8xi32, #blocked>
+      %71 = tt.addptr %47, %70 : tensor<64x8x!tt.ptr<bf16, 1>, #blocked>, tensor<64x8xi32, #blocked>
+      %72 = arith.truncf %69 : tensor<64x8xf32, #blocked> to tensor<64x8xbf16, #blocked>
+      tt.store %71, %72, %54 {cache = 1 : i32, evict = 1 : i32} : tensor<64x8xbf16, #blocked>
+    }
+    tt.return
+  }
+}

.triton/dump/d7a12c0ba96f8920b8147157303ee99f/triton_.ptx

@@ -0,0 +1,723 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2d3d4d5de6de
+.extern .shared .align 1 .b8 global_smem[];
+.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};
+
+.visible .entry triton__0d1d2d3d4d5de6de(
+	.param .u64 triton__0d1d2d3d4d5de6de_param_0,
+	.param .u64 triton__0d1d2d3d4d5de6de_param_1,
+	.param .u64 triton__0d1d2d3d4d5de6de_param_2,
+	.param .u64 triton__0d1d2d3d4d5de6de_param_3,
+	.param .u64 triton__0d1d2d3d4d5de6de_param_4,
+	.param .u32 triton__0d1d2d3d4d5de6de_param_5,
+	.param .u32 triton__0d1d2d3d4d5de6de_param_6
+)
+.maxntid 64, 1, 1
+{
+	.reg .pred 	%p<26>;
+	.reg .b16 	%rs<13>;
+	.reg .b32 	%r<92>;
+	.reg .f32 	%f<78>;
+	.reg .b64 	%rd<14>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd6, [triton__0d1d2d3d4d5de6de_param_0];
+	ld.param.u64 	%rd7, [triton__0d1d2d3d4d5de6de_param_1];
+$L__tmp0:
+	.loc	1 26 26
+	mov.u32 	%r58, %tid.x;
+	and.b32  	%r59, %r58, 31;
+	ld.param.u64 	%rd8, [triton__0d1d2d3d4d5de6de_param_2];
+	ld.param.u64 	%rd9, [triton__0d1d2d3d4d5de6de_param_3];
+	ld.param.u64 	%rd10, [triton__0d1d2d3d4d5de6de_param_4];
+	shl.b32 	%r60, %r58, 2;
+	and.b32  	%r61, %r60, 252;
+	.loc	1 23 28
+	mov.u32 %r1, %ctaid.x;
+	.loc	1 30 40
+	shl.b32 	%r62, %r1, 8;
+	.loc	1 30 36
+	or.b32  	%r63, %r62, %r61;
+	.loc	1 30 30
+	mul.wide.s32 	%rd11, %r63, 4;
+	add.s64 	%rd1, %rd6, %rd11;
+	mov.b32 	%r6, 0;
+	mov.pred 	%p1, -1;
+	.loc	1 30 46
+	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 ];
+	@!%p1 mov.u32 %r2, %r6;
+	@!%p1 mov.u32 %r3, %r6;
+	@!%p1 mov.u32 %r4, %r6;
+	@!%p1 mov.u32 %r5, %r6;
+	mov.b32 	%f1, %r4;
+	mov.b32 	%f2, %r5;
+	.loc	1 31 30
+	mul.wide.s32 	%rd12, %r63, 2;
+	add.s64 	%rd2, %rd7, %rd12;
+	.loc	1 31 46
+	mov.u32 %r10, 0x0;
+	mov.u32 %r11, 0x0;
+	@%p1 ld.global.v2.b32 { %r10, %r11 }, [ %rd2 + 0 ];
+	@!%p1 mov.u32 %r10, %r6;
+	@!%p1 mov.u32 %r11, %r6;
+	cvt.u16.u32 	%rs1, %r10;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r10; }
+	cvt.u16.u32 	%rs3, %r11;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r11; }
+	.loc	1 31 67
+	cvt.f32.bf16 %r14, %rs1;
+	mov.b32 	%f3, %r14;
+	cvt.f32.bf16 %r15, %rs2;
+	mov.b32 	%f4, %r15;
+	cvt.f32.bf16 %r16, %rs3;
+	mov.b32 	%f5, %r16;
+	cvt.f32.bf16 %r17, %rs4;
+	mov.b32 	%f6, %r17;
+	.loc	1 32 30
+	add.s64 	%rd3, %rd8, %rd12;
+	.loc	1 32 46
+	mov.u32 %r18, 0x0;
+	mov.u32 %r19, 0x0;
+	@%p1 ld.global.v2.b32 { %r18, %r19 }, [ %rd3 + 0 ];
+	@!%p1 mov.u32 %r18, %r6;
+	@!%p1 mov.u32 %r19, %r6;
+	cvt.u16.u32 	%rs5, %r18;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r18; }
+	cvt.u16.u32 	%rs7, %r19;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r19; }
+	.loc	1 32 67
+	cvt.f32.bf16 %r22, %rs5;
+	mov.b32 	%f7, %r22;
+	cvt.f32.bf16 %r23, %rs6;
+	mov.b32 	%f8, %r23;
+	cvt.f32.bf16 %r24, %rs7;
+	mov.b32 	%f9, %r24;
+	cvt.f32.bf16 %r25, %rs8;
+	mov.b32 	%f10, %r25;
+	.loc	1 33 31
+	mul.wide.u32 	%rd13, %r61, 4;
+	add.s64 	%rd4, %rd9, %rd13;
+	.loc	1 33 36
+	mov.u32 %r26, 0x0;
+	mov.u32 %r27, 0x0;
+	mov.u32 %r28, 0x0;
+	mov.u32 %r29, 0x0;
+	@%p1 ld.global.L1::evict_last.v4.b32 { %r26, %r27, %r28, %r29 }, [ %rd4 + 0 ];
+	@!%p1 mov.u32 %r26, %r6;
+	@!%p1 mov.u32 %r27, %r6;
+	@!%p1 mov.u32 %r28, %r6;
+	@!%p1 mov.u32 %r29, %r6;
+	.loc	1 35 18
+	add.f32 	%f11, %f5, %f1;
+	add.f32 	%f12, %f6, %f2;
+	.loc	1 30 46
+	mov.b32 	%f13, %r3;
+	mov.b32 	%f14, %r2;
+	.loc	1 35 18
+	add.f32 	%f15, %f3, %f14;
+	add.f32 	%f16, %f4, %f13;
+	.loc	1 37 18
+	add.f32 	%f17, %f16, %f8;
+	add.f32 	%f18, %f15, %f7;
+	add.f32 	%f19, %f11, %f9;
+	add.f32 	%f20, %f12, %f10;
+$L__tmp1:
+	.loc	2 233 15
+	add.f32 	%f21, %f18, %f17;
+	add.f32 	%f22, %f21, %f19;
+	add.f32 	%f23, %f22, %f20;
+$L__tmp2:
+	.loc	2 243 36
+	mov.b32 	%r64, %f23;
+	shfl.sync.bfly.b32	%r65, %r64, 16, 31, -1;
+	mov.b32 	%f24, %r65;
+$L__tmp3:
+	.loc	2 233 15
+	add.f32 	%f25, %f23, %f24;
+$L__tmp4:
+	.loc	2 243 36
+	mov.b32 	%r66, %f25;
+	shfl.sync.bfly.b32	%r67, %r66, 8, 31, -1;
+	mov.b32 	%f26, %r67;
+$L__tmp5:
+	.loc	2 233 15
+	add.f32 	%f27, %f25, %f26;
+$L__tmp6:
+	.loc	2 243 36
+	mov.b32 	%r68, %f27;
+	shfl.sync.bfly.b32	%r69, %r68, 4, 31, -1;
+	mov.b32 	%f28, %r69;
+$L__tmp7:
+	.loc	2 233 15
+	add.f32 	%f29, %f27, %f28;
+$L__tmp8:
+	.loc	2 243 36
+	mov.b32 	%r70, %f29;
+	shfl.sync.bfly.b32	%r71, %r70, 2, 31, -1;
+	mov.b32 	%f30, %r71;
+$L__tmp9:
+	.loc	2 233 15
+	add.f32 	%f31, %f29, %f30;
+$L__tmp10:
+	.loc	2 243 36
+	mov.b32 	%r72, %f31;
+	shfl.sync.bfly.b32	%r73, %r72, 1, 31, -1;
+	mov.b32 	%f32, %r73;
+$L__tmp11:
+	.loc	2 233 15
+	add.f32 	%f33, %f31, %f32;
+$L__tmp12:
+	.loc	2 243 36
+	setp.eq.s32 	%p17, %r59, 0;
+	shr.u32 	%r74, %r58, 3;
+	and.b32  	%r75, %r74, 4;
+	mov.u32 	%r76, global_smem;
+	add.s32 	%r34, %r76, %r75;
+	mov.b32 	%r35, %f33;
+	@%p17 st.shared.b32 [ %r34 + 0 ], %r35;
+	bar.sync 	0;
+	setp.lt.s32 	%p18, %r58, 2;
+	add.s32 	%r37, %r76, %r60;
+	@%p18 ld.shared.b32 %r36, [ %r37 + 0 ];
+	mov.b32 	%f34, %r36;
+	shfl.sync.bfly.b32	%r77, %r36, 1, 31, -1;
+	mov.b32 	%f35, %r77;
+$L__tmp13:
+	.loc	2 233 15
+	add.f32 	%f36, %f34, %f35;
+$L__tmp14:
+	.loc	2 243 36
+	and.b32  	%r78, %r58, 1;
+	setp.eq.b32 	%p24, %r78, 1;
+	not.pred 	%p25, %p24;
+	and.pred  	%p19, %p18, %p25;
+	mov.b32 	%r39, %f36;
+	@%p19 st.shared.b32 [ %r37 + 0 ], %r39;
+	bar.sync 	0;
+	ld.shared.f32 	%f37, [global_smem];
+$L__tmp15:
+	.loc	3 8 15
+	add.f32 	%f38, %f37, 0f00000000;
+$L__tmp16:
+	.loc	1 45 20
+	mov.b32 	%r41, %f38;
+	mov.b32 	%r42, 1132462080;
+	div.full.f32 %r40, %r41, %r42;
+	mov.b32 	%f39, %r40;
+	.loc	1 46 19
+	sub.f32 	%f40, %f18, %f39;
+	sub.f32 	%f41, %f17, %f39;
+	sub.f32 	%f42, %f19, %f39;
+	sub.f32 	%f43, %f20, %f39;
+	.loc	1 47 20
+	mul.f32 	%f44, %f41, %f41;
+$L__tmp17:
+	.loc	2 243 36
+	bar.sync 	0;
+$L__tmp18:
+	.loc	2 233 15
+	fma.rn.f32 	%f45, %f40, %f40, %f44;
+	fma.rn.f32 	%f46, %f42, %f42, %f45;
+	fma.rn.f32 	%f47, %f43, %f43, %f46;
+$L__tmp19:
+	.loc	2 243 36
+	mov.b32 	%r79, %f47;
+	shfl.sync.bfly.b32	%r80, %r79, 16, 31, -1;
+	mov.b32 	%f48, %r80;
+$L__tmp20:
+	.loc	2 233 15
+	add.f32 	%f49, %f47, %f48;
+$L__tmp21:
+	.loc	2 243 36
+	mov.b32 	%r81, %f49;
+	shfl.sync.bfly.b32	%r82, %r81, 8, 31, -1;
+	mov.b32 	%f50, %r82;
+$L__tmp22:
+	.loc	2 233 15
+	add.f32 	%f51, %f49, %f50;
+$L__tmp23:
+	.loc	2 243 36
+	mov.b32 	%r83, %f51;
+	shfl.sync.bfly.b32	%r84, %r83, 4, 31, -1;
+	mov.b32 	%f52, %r84;
+$L__tmp24:
+	.loc	2 233 15
+	add.f32 	%f53, %f51, %f52;
+$L__tmp25:
+	.loc	2 243 36
+	mov.b32 	%r85, %f53;
+	shfl.sync.bfly.b32	%r86, %r85, 2, 31, -1;
+	mov.b32 	%f54, %r86;
+$L__tmp26:
+	.loc	2 233 15
+	add.f32 	%f55, %f53, %f54;
+$L__tmp27:
+	.loc	2 243 36
+	mov.b32 	%r87, %f55;
+	shfl.sync.bfly.b32	%r88, %r87, 1, 31, -1;
+	mov.b32 	%f56, %r88;
+$L__tmp28:
+	.loc	2 233 15
+	add.f32 	%f57, %f55, %f56;
+$L__tmp29:
+	.loc	2 243 36
+	mov.b32 	%r44, %f57;
+	@%p17 st.shared.b32 [ %r34 + 0 ], %r44;
+	bar.sync 	0;
+	@%p18 ld.shared.b32 %r45, [ %r37 + 0 ];
+	mov.b32 	%f58, %r45;
+	shfl.sync.bfly.b32	%r89, %r45, 1, 31, -1;
+	mov.b32 	%f59, %r89;
+$L__tmp30:
+	.loc	2 233 15
+	add.f32 	%f60, %f58, %f59;
+$L__tmp31:
+	.loc	2 243 36
+	mov.b32 	%r48, %f60;
+	@%p19 st.shared.b32 [ %r37 + 0 ], %r48;
+	bar.sync 	0;
+	ld.shared.f32 	%f61, [global_smem];
+$L__tmp32:
+	.loc	3 8 15
+	add.f32 	%f62, %f61, 0f00000000;
+$L__tmp33:
+	.loc	1 53 20
+	mov.b32 	%r50, %f62;
+	div.full.f32 %r49, %r50, %r42;
+	mov.b32 	%f63, %r49;
+	.loc	1 55 20
+	add.f32 	%f64, %f63, 0f3727C5AC;
+	.loc	1 56 26
+	rsqrt.approx.ftz.f32 	%f65, %f64;
+	.loc	1 33 36
+	mov.b32 	%f66, %r29;
+	mov.b32 	%f67, %r28;
+	mov.b32 	%f68, %r27;
+	mov.b32 	%f69, %r26;
+	.loc	1 57 20
+	mul.f32 	%f70, %f40, %f65;
+	mul.f32 	%f71, %f41, %f65;
+	mul.f32 	%f72, %f42, %f65;
+	mul.f32 	%f73, %f43, %f65;
+	.loc	1 58 20
+	mul.f32 	%f74, %f70, %f69;
+	mul.f32 	%f75, %f71, %f68;
+	mul.f32 	%f76, %f72, %f67;
+	mul.f32 	%f77, %f73, %f66;
+	.loc	1 60 25
+	add.s64 	%rd5, %rd10, %rd12;
+	.loc	1 60 48
+	mov.b32 	%r52, %f74;
+	cvt.rn.bf16.f32 %rs9, %r52;
+	mov.b32 	%r53, %f75;
+	cvt.rn.bf16.f32 %rs10, %r53;
+	mov.b32 	%r54, %f76;
+	cvt.rn.bf16.f32 %rs11, %r54;
+	mov.b32 	%r55, %f77;
+	cvt.rn.bf16.f32 %rs12, %r55;
+	mov.b32 	%r90, {%rs9, %rs10};
+	mov.b32 	%r91, {%rs11, %rs12};
+	@%p1 st.global.v2.b32 [ %rd5 + 0 ], { %r90, %r91 };
+	.loc	1 60 4
+	ret;
+$L__tmp34:
+$L__func_end0:
+
+}
+	// .globl	__nv_rsqrtf
+.visible .func  (.param .b32 func_retval0) __nv_rsqrtf(
+	.param .b32 __nv_rsqrtf_param_0
+)
+{
+	.reg .f32 	%f<3>;
+$L__func_begin1:
+
+	ld.param.f32 	%f1, [__nv_rsqrtf_param_0];
+	rsqrt.approx.ftz.f32 	%f2, %f1;
+	st.param.f32 	[func_retval0+0], %f2;
+	ret;
+$L__func_end1:
+
+}
+	.file	1 "/tmp/torchinductor_root/do/cdohrmmhfsykzlva6pepxaa7gf7klw7w5jzorpspyaldhfg3acr2.py"
+	.file	2 "/usr/local/lib/python3.10/dist-packages/triton/language/standard.py"
+	.file	3 "/usr/local/lib/python3.10/dist-packages/torch/_inductor/triton_helpers.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 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 5
+.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 395
+.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 100
+.b8 111
+.b8 104
+.b8 114
+.b8 109
+.b8 109
+.b8 104
+.b8 102
+.b8 115
+.b8 121
+.b8 107
+.b8 122
+.b8 108
+.b8 118
+.b8 97
+.b8 54
+.b8 112
+.b8 101
+.b8 112
+.b8 120
+.b8 97
+.b8 97
+.b8 55
+.b8 103
+.b8 102
+.b8 55
+.b8 107
+.b8 108
+.b8 119
+.b8 55
+.b8 119
+.b8 53
+.b8 106
+.b8 122
+.b8 111
+.b8 114
+.b8 112
+.b8 115
+.b8 112
+.b8 121
+.b8 97
+.b8 108
+.b8 100
+.b8 104
+.b8 102
+.b8 103
+.b8 51
+.b8 97
+.b8 99
+.b8 114
+.b8 50
+.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 100
+.b8 111
+.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 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 101
+.b8 54
+.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 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 101
+.b8 54
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 1
+.b8 3
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b32 125
+.b8 4
+.b32 125
+.b64 $L__tmp1
+.b64 $L__tmp14
+.b8 2
+.b8 42
+.b8 59
+.b8 5
+.b32 125
+.b64 $L__tmp1
+.b64 $L__tmp14
+.b8 2
+.b8 243
+.b8 36
+.b8 0
+.b8 5
+.b32 125
+.b64 $L__tmp2
+.b64 $L__tmp15
+.b8 2
+.b8 42
+.b8 59
+.b8 5
+.b32 125
+.b64 $L__tmp15
+.b64 $L__tmp16
+.b8 3
+.b8 42
+.b8 45
+.b8 5
+.b32 125
+.b64 $L__tmp17
+.b64 $L__tmp32
+.b8 2
+.b8 50
+.b8 59
+.b8 4
+.b32 125
+.b64 $L__tmp18
+.b64 $L__tmp31
+.b8 2
+.b8 50
+.b8 59
+.b8 5
+.b32 125
+.b64 $L__tmp18
+.b64 $L__tmp31
+.b8 2
+.b8 243
+.b8 36
+.b8 0
+.b8 5
+.b32 125
+.b64 $L__tmp32
+.b64 $L__tmp33
+.b8 3
+.b8 50
+.b8 45
+.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 399
+.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 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 101
+.b8 54
+.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 399
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}