Add files using upload-large-folder tool

SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/__grp__triton_red_fused__to_copy_clone_slice_sum_transpose_5.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_red_fused__to_copy_clone_slice_sum_transpose_5.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.source", "triton_red_fused__to_copy_clone_slice_sum_transpose_5.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.ttir", "triton_red_fused__to_copy_clone_slice_sum_transpose_5.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.ttgir", "triton_red_fused__to_copy_clone_slice_sum_transpose_5.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.llir", "triton_red_fused__to_copy_clone_slice_sum_transpose_5.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.ptx", "triton_red_fused__to_copy_clone_slice_sum_transpose_5.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.cubin", "triton_red_fused__to_copy_clone_slice_sum_transpose_5.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.json"}}

SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.json

@@ -0,0 +1 @@
+{"hash": "d1c2e6527ce27b628a96c1a250025b39aad19d679fce295e820390aa7ae64b66", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 4, "num_ctas": 1, "num_stages": 1, "warp_size": 32, "maxnreg": null, "cluster_dims": [1, 1, 1], "ptx_version": null, "ptx_options": null, "ir_override": null, "enable_fp_fusion": true, "launch_cooperative_grid": false, "launch_pdl": false, "supported_fp8_dtypes": ["fp8e4b15", "fp8e4nv", "fp8e5"], "deprecated_fp8_dot_operand_dtypes": ["fp8e4b15"], "default_dot_input_precision": "tf32", "allowed_dot_input_precisions": ["tf32", "tf32x3", "ieee"], "max_num_imprecise_acc_default": 1073741824, "extern_libs": [["libdevice", "/workspace/specforge/lib/python3.11/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]], "debug": true, "backend_name": "cuda", "sanitize_overflow": false, "arch": "sm90", "instrumentation_mode": "", "triton_version": "3.5.1", "tensordesc_meta": [], "shared": 1024, "tmem_size": 0, "global_scratch_size": 0, "global_scratch_align": 1, "profile_scratch_size": 0, "profile_scratch_align": 1, "name": "triton_red_fused__to_copy_clone_slice_sum_transpose_5"}

SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.ptx

@@ -0,0 +1,525 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.7
+.target sm_90a
+.address_size 64
+
+	// .globl	triton_red_fused__to_copy_clone_slice_sum_transpose_5 // -- Begin function triton_red_fused__to_copy_clone_slice_sum_transpose_5
+.extern .shared .align 16 .b8 global_smem[];
+                                        // @triton_red_fused__to_copy_clone_slice_sum_transpose_5
+.visible .entry triton_red_fused__to_copy_clone_slice_sum_transpose_5(
+	.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_0,
+	.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_1,
+	.param .u64 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_2,
+	.param .u64 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_3,
+	.param .u32 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_4,
+	.param .u32 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_5,
+	.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_6,
+	.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_7
+)
+.reqntid 128
+{
+	.reg .pred 	%p<24>;
+	.reg .b32 	%r<51>;
+	.reg .b64 	%rd<97>;
+	.loc	1 18 0                          // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:18:0
+$L__func_begin0:
+	.loc	1 18 0                          // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:18:0
+
+// %bb.0:
+	ld.param.b32 	%r11, [triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_5];
+	ld.param.b64 	%rd20, [triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_2];
+$L__tmp0:
+	.loc	1 21 28                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:21:28
+	mov.u32 	%r12, %ctaid.x;
+	.loc	1 21 33                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:21:33
+	shl.b32 	%r13, %r12, 5;
+	.loc	1 22 44                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:22:44
+	mov.u32 	%r1, %tid.x;
+	and.b32 	%r2, %r1, 31;
+	.loc	1 22 23                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:22:23
+	or.b32 	%r14, %r13, %r2;
+	.loc	1 26 19                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:26:19
+	cvt.s64.s32 	%rd1, %r14;
+	.loc	1 27 19                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:27:19
+	or.b64 	%rd23, %rd1, %rd20;
+	and.b64 	%rd24, %rd23, -4294967296;
+	setp.ne.b64 	%p5, %rd24, 0;
+	cvt.u32.u64 	%r49, %rd1;
+	@%p5 bra 	$L__BB0_2;
+	bra.uni 	$L__BB0_1;
+$L__BB0_2:
+	div.s64 	%rd91, %rd1, %rd20;
+	bra.uni 	$L__BB0_3;
+$L__BB0_1:
+	cvt.u32.u64 	%r15, %rd20;
+	div.u32 	%r17, %r49, %r15;
+	cvt.u64.u32 	%rd91, %r17;
+$L__BB0_3:
+	.loc	1 0 19                          // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:0:19
+	ld.param.b32 	%r10, [triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_4];
+	ld.param.b64 	%rd19, [triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_1];
+	bfe.u32 	%r3, %r1, 5, 2;
+	.loc	1 26 19                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:26:19
+	mul.lo.s64 	%rd26, %rd91, %rd20;
+	sub.s64 	%rd6, %rd1, %rd26;
+	.loc	1 30 40                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:30:40
+	setp.lt.s32 	%p6, %r11, 1;
+	mov.b64 	%rd96, 0;
+	shl.b64 	%rd90, %rd6, 2;
+	@%p6 bra 	$L__BB0_7;
+// %bb.4:                               // %.lr.ph
+	.loc	1 0 40                          // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:0:40
+	ld.param.b64 	%rd21, [triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_3];
+	ld.param.b64 	%rd18, [triton_red_fused__to_copy_clone_slice_sum_transpose_5_param_0];
+	.loc	1 23 21                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:23:21
+	setp.lt.s32 	%p1, %r49, %r10;
+	.loc	1 36 54                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:54
+	mul.lo.s64 	%rd31, %rd21, %rd20;
+	.loc	1 36 58                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:58
+	mul.lo.s64 	%rd32, %rd31, %rd91;
+	add.s64 	%rd34, %rd18, %rd90;
+	.loc	1 30 40                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:30:40
+	shl.b64 	%rd35, %rd32, 2;
+	add.s64 	%rd7, %rd34, %rd35;
+	mov.b64 	%rd92, 0;
+	mov.b32 	%r50, 0;
+	mov.b64 	%rd93, %rd92;
+	mov.b64 	%rd94, %rd92;
+	mov.b64 	%rd95, %rd92;
+$L__BB0_5:                              // =>This Inner Loop Header: Depth=1
+	.loc	1 36 63                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:63
+	// begin inline asm
+	mov.u64 %rd36, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd36, 1.0;
+	// end inline asm
+	.loc	1 31 31                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:31:31
+	add.s32 	%r24, %r3, %r50;
+	add.s32 	%r25, %r24, 4;
+	add.s32 	%r26, %r24, 8;
+	.loc	1 32 29                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:32:29
+	add.s32 	%r27, %r24, 12;
+	setp.lt.s32 	%p11, %r24, %r11;
+	setp.lt.s32 	%p12, %r25, %r11;
+	setp.lt.s32 	%p13, %r26, %r11;
+	setp.lt.s32 	%p14, %r27, %r11;
+	.loc	1 36 43                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:43
+	cvt.s64.s32 	%rd48, %r24;
+	cvt.s64.s32 	%rd49, %r25;
+	cvt.s64.s32 	%rd50, %r26;
+	cvt.s64.s32 	%rd51, %r27;
+	mul.lo.s64 	%rd52, %rd20, %rd48;
+	mul.lo.s64 	%rd53, %rd20, %rd49;
+	mul.lo.s64 	%rd54, %rd20, %rd50;
+	mul.lo.s64 	%rd55, %rd20, %rd51;
+	.loc	1 36 34                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:34
+	shl.b64 	%rd56, %rd52, 2;
+	add.s64 	%rd37, %rd7, %rd56;
+	shl.b64 	%rd57, %rd53, 2;
+	add.s64 	%rd40, %rd7, %rd57;
+	shl.b64 	%rd58, %rd54, 2;
+	add.s64 	%rd43, %rd7, %rd58;
+	shl.b64 	%rd59, %rd55, 2;
+	add.s64 	%rd46, %rd7, %rd59;
+	.loc	1 36 73                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:73
+	and.pred 	%p10, %p1, %p14;
+	and.pred 	%p9, %p1, %p13;
+	and.pred 	%p8, %p1, %p12;
+	and.pred 	%p7, %p1, %p11;
+	.loc	1 36 63                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:36:63
+	// begin inline asm
+	mov.u32 %r20, 0x0;
+	@%p7 ld.global.L1::evict_last.L2::cache_hint.b32 { %r20 }, [ %rd37 + 0 ], %rd36;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd39, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd39, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r21, 0x0;
+	@%p8 ld.global.L1::evict_last.L2::cache_hint.b32 { %r21 }, [ %rd40 + 0 ], %rd39;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd42, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd42, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r22, 0x0;
+	@%p9 ld.global.L1::evict_last.L2::cache_hint.b32 { %r22 }, [ %rd43 + 0 ], %rd42;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd45, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd45, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r23, 0x0;
+	@%p10 ld.global.L1::evict_last.L2::cache_hint.b32 { %r23 }, [ %rd46 + 0 ], %rd45;
+	// end inline asm
+	.loc	1 37 23                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:37:23
+	cvt.s64.s32 	%rd60, %r20;
+	cvt.s64.s32 	%rd61, %r21;
+	cvt.s64.s32 	%rd62, %r22;
+	cvt.s64.s32 	%rd63, %r23;
+	.loc	1 40 48                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:40:48
+	selp.b64 	%rd64, %rd63, 0, %p10;
+	selp.b64 	%rd65, %rd62, 0, %p9;
+	selp.b64 	%rd66, %rd61, 0, %p8;
+	selp.b64 	%rd67, %rd60, 0, %p7;
+	add.s64 	%rd92, %rd67, %rd92;
+	add.s64 	%rd93, %rd66, %rd93;
+	add.s64 	%rd94, %rd65, %rd94;
+	add.s64 	%rd95, %rd64, %rd95;
+	.loc	1 30 40                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:30:40
+	add.s32 	%r50, %r50, 16;
+	setp.lt.s32 	%p15, %r50, %r11;
+	@%p15 bra 	$L__BB0_5;
+// %bb.6:                               // %._crit_edge.loopexit
+$L__tmp1:
+	.loc	2 261 15                        // standard.py:261:15 @[ c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:41:25 ]
+	add.s64 	%rd68, %rd92, %rd94;
+	add.s64 	%rd69, %rd93, %rd95;
+	add.s64 	%rd96, %rd68, %rd69;
+$L__tmp2:
+$L__BB0_7:                              // %._crit_edge
+	.loc	1 23 21                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:23:21
+	setp.lt.s32 	%p20, %r49, %r10;
+$L__tmp3:
+	.loc	2 291 36                        // standard.py:291:36 @[ c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:41:25 ]
+	shl.b32 	%r33, %r2, 5;
+	mov.b32 	%r34, global_smem;
+	add.s32 	%r35, %r34, %r33;
+	shl.b32 	%r36, %r3, 3;
+	add.s32 	%r28, %r35, %r36;
+	mov.pred 	%p16, -1;
+	// begin inline asm
+	@%p16 st.shared.b64 [ %r28 + 0 ], %rd96;
+	// end inline asm
+	bar.sync 	0;
+	setp.lt.u32 	%p17, %r1, 128;
+	shl.b32 	%r37, %r1, 3;
+	add.s32 	%r29, %r34, %r37;
+	// begin inline asm
+	@%p17 ld.shared.b64 %rd71, [ %r29 + 0 ];
+	// end inline asm
+	mov.b64 	{_, %r38}, %rd71;
+	cvt.u32.u64 	%r39, %rd71;
+	shfl.sync.bfly.b32 	%r40, %r39, 2, 31, -1;
+	shfl.sync.bfly.b32 	%r41, %r38, 2, 31, -1;
+	cvt.u64.u32 	%rd74, %r40;
+	cvt.u64.u32 	%rd75, %r41;
+	shl.b64 	%rd76, %rd75, 32;
+	or.b64 	%rd77, %rd74, %rd76;
+	.loc	2 261 15                        // standard.py:261:15 @[ c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:41:25 ]
+	add.s64 	%rd78, %rd71, %rd77;
+	.loc	2 291 36                        // standard.py:291:36 @[ c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:41:25 ]
+	mov.b64 	{_, %r42}, %rd78;
+	cvt.u32.u64 	%r43, %rd78;
+	shfl.sync.bfly.b32 	%r44, %r43, 1, 31, -1;
+	shfl.sync.bfly.b32 	%r45, %r42, 1, 31, -1;
+	cvt.u64.u32 	%rd79, %r44;
+	cvt.u64.u32 	%rd80, %r45;
+	shl.b64 	%rd81, %rd80, 32;
+	or.b64 	%rd82, %rd79, %rd81;
+	.loc	2 261 15                        // standard.py:261:15 @[ c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:41:25 ]
+	add.s64 	%rd72, %rd78, %rd82;
+	.loc	2 291 36                        // standard.py:291:36 @[ c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:41:25 ]
+	and.b32 	%r46, %r1, 899;
+	setp.eq.b32 	%p18, %r46, 0;
+	// begin inline asm
+	@%p18 st.shared.b64 [ %r29 + 0 ], %rd72;
+	// end inline asm
+	bar.sync 	0;
+	ld.shared.b32 	%r31, [%r35];
+$L__tmp4:
+	.loc	1 43 49                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:49
+	setp.lt.s64 	%p21, %rd20, 2;
+	.loc	1 43 75                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:75
+	setp.gt.s64 	%p22, %rd20, 1;
+	.loc	1 43 66                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:66
+	selp.b64 	%rd83, %rd20, 0, %p22;
+	.loc	1 43 0                          // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43
+	selp.b64 	%rd84, 1, 0, %p21;
+	.loc	1 43 57                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:57
+	add.s64 	%rd85, %rd83, %rd84;
+	.loc	1 43 34                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:34
+	mul.lo.s64 	%rd86, %rd91, %rd85;
+	.loc	1 43 25                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:25
+	add.s64 	%rd88, %rd19, %rd90;
+	shl.b64 	%rd89, %rd86, 2;
+	add.s64 	%rd73, %rd88, %rd89;
+	.loc	1 43 88                         // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:88
+	and.b32 	%r47, %r1, 96;
+	setp.eq.b32 	%p23, %r47, 0;
+	and.pred 	%p19, %p23, %p20;
+	// begin inline asm
+	@%p19 st.global.b32 [ %rd73 + 0 ], { %r31 };
+	// end inline asm
+	.loc	1 43 4                          // c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py:43:4
+	ret;
+$L__tmp5:
+$L__func_end0:
+                                        // -- End function
+}
+	.file	1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py"
+	.file	2 "/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py"
+	.section	.debug_abbrev
+	{
+.b8 1                                   // Abbreviation Code
+.b8 17                                  // DW_TAG_compile_unit
+.b8 1                                   // DW_CHILDREN_yes
+.b8 37                                  // DW_AT_producer
+.b8 8                                   // DW_FORM_string
+.b8 19                                  // DW_AT_language
+.b8 5                                   // DW_FORM_data2
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 16                                  // DW_AT_stmt_list
+.b8 6                                   // DW_FORM_data4
+.b8 27                                  // DW_AT_comp_dir
+.b8 8                                   // DW_FORM_string
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 2                                   // Abbreviation Code
+.b8 46                                  // DW_TAG_subprogram
+.b8 0                                   // DW_CHILDREN_no
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 32                                  // DW_AT_inline
+.b8 11                                  // DW_FORM_data1
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 3                                   // Abbreviation Code
+.b8 46                                  // DW_TAG_subprogram
+.b8 1                                   // DW_CHILDREN_yes
+.b8 17                                  // DW_AT_low_pc
+.b8 1                                   // DW_FORM_addr
+.b8 18                                  // DW_AT_high_pc
+.b8 1                                   // DW_FORM_addr
+.b8 49                                  // DW_AT_abstract_origin
+.b8 19                                  // DW_FORM_ref4
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 4                                   // Abbreviation Code
+.b8 29                                  // DW_TAG_inlined_subroutine
+.b8 0                                   // DW_CHILDREN_no
+.b8 49                                  // DW_AT_abstract_origin
+.b8 19                                  // DW_FORM_ref4
+.b8 17                                  // DW_AT_low_pc
+.b8 1                                   // DW_FORM_addr
+.b8 18                                  // DW_AT_high_pc
+.b8 1                                   // DW_FORM_addr
+.b8 88                                  // DW_AT_call_file
+.b8 11                                  // DW_FORM_data1
+.b8 89                                  // DW_AT_call_line
+.b8 11                                  // DW_FORM_data1
+.b8 87                                  // DW_AT_call_column
+.b8 11                                  // DW_FORM_data1
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 0                                   // EOM(3)
+	}
+	.section	.debug_info
+	{
+.b32 238                                // Length of Unit
+.b8 2                                   // DWARF version number
+.b8 0
+.b32 .debug_abbrev                      // Offset Into Abbrev. Section
+.b8 8                                   // Address Size (in bytes)
+.b8 1                                   // Abbrev [1] 0xb:0xe7 DW_TAG_compile_unit
+.b8 116                                 // DW_AT_producer
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2                                   // DW_AT_language
+.b8 0
+.b8 99                                  // DW_AT_name
+.b8 53
+.b8 117
+.b8 99
+.b8 102
+.b8 50
+.b8 105
+.b8 116
+.b8 54
+.b8 102
+.b8 119
+.b8 104
+.b8 109
+.b8 114
+.b8 117
+.b8 51
+.b8 50
+.b8 118
+.b8 112
+.b8 121
+.b8 50
+.b8 120
+.b8 109
+.b8 119
+.b8 107
+.b8 107
+.b8 52
+.b8 118
+.b8 103
+.b8 55
+.b8 112
+.b8 112
+.b8 110
+.b8 111
+.b8 120
+.b8 109
+.b8 104
+.b8 110
+.b8 107
+.b8 102
+.b8 54
+.b8 105
+.b8 102
+.b8 112
+.b8 98
+.b8 50
+.b8 104
+.b8 50
+.b8 116
+.b8 106
+.b8 115
+.b8 108
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line                        // DW_AT_stmt_list
+.b8 47                                  // DW_AT_comp_dir
+.b8 119
+.b8 111
+.b8 114
+.b8 107
+.b8 115
+.b8 112
+.b8 97
+.b8 99
+.b8 101
+.b8 47
+.b8 104
+.b8 97
+.b8 110
+.b8 114
+.b8 117
+.b8 105
+.b8 47
+.b8 83
+.b8 112
+.b8 101
+.b8 99
+.b8 70
+.b8 111
+.b8 114
+.b8 103
+.b8 101
+.b8 45
+.b8 101
+.b8 120
+.b8 116
+.b8 47
+.b8 99
+.b8 97
+.b8 99
+.b8 104
+.b8 101
+.b8 47
+.b8 99
+.b8 111
+.b8 109
+.b8 112
+.b8 105
+.b8 108
+.b8 101
+.b8 100
+.b8 95
+.b8 107
+.b8 101
+.b8 114
+.b8 110
+.b8 101
+.b8 108
+.b8 115
+.b8 47
+.b8 53
+.b8 117
+.b8 0
+.b8 2                                   // Abbrev [2] 0x8b:0x38 DW_TAG_subprogram
+.b8 116                                 // DW_AT_name
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 114
+.b8 101
+.b8 100
+.b8 95
+.b8 102
+.b8 117
+.b8 115
+.b8 101
+.b8 100
+.b8 95
+.b8 95
+.b8 116
+.b8 111
+.b8 95
+.b8 99
+.b8 111
+.b8 112
+.b8 121
+.b8 95
+.b8 99
+.b8 108
+.b8 111
+.b8 110
+.b8 101
+.b8 95
+.b8 115
+.b8 108
+.b8 105
+.b8 99
+.b8 101
+.b8 95
+.b8 115
+.b8 117
+.b8 109
+.b8 95
+.b8 116
+.b8 114
+.b8 97
+.b8 110
+.b8 115
+.b8 112
+.b8 111
+.b8 115
+.b8 101
+.b8 95
+.b8 53
+.b8 0
+.b8 1                                   // DW_AT_inline
+.b8 3                                   // Abbrev [3] 0xc3:0x2e DW_TAG_subprogram
+.b64 $L__func_begin0                    // DW_AT_low_pc
+.b64 $L__func_end0                      // DW_AT_high_pc
+.b32 139                                // DW_AT_abstract_origin
+.b8 4                                   // Abbrev [4] 0xd8:0x18 DW_TAG_inlined_subroutine
+.b32 139                                // DW_AT_abstract_origin
+.b64 $L__tmp1                           // DW_AT_low_pc
+.b64 $L__tmp4                           // DW_AT_high_pc
+.b8 1                                   // DW_AT_call_file
+.b8 41                                  // DW_AT_call_line
+.b8 25                                  // DW_AT_call_column
+.b8 0                                   // End Of Children Mark
+.b8 0                                   // End Of Children Mark
+	}
+	.section	.debug_macinfo	{	}

SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.ttgir

@@ -0,0 +1,147 @@
+#blocked = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [1, 4], order = [0, 1]}>
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":18:0)
+#loc1 = loc(unknown)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":41:25)
+#loc40 = loc("in_ptr0"(#loc))
+#loc41 = loc("out_ptr1"(#loc))
+#loc42 = loc("ks0"(#loc))
+#loc43 = loc("ks1"(#loc))
+#loc44 = loc("xnumel"(#loc))
+#loc45 = loc("r0_numel"(#loc))
+#loc68 = loc("tmp3"(#loc26))
+#loc73 = loc(callsite(#loc1 at #loc68))
+module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
+  tt.func public @triton_red_fused__to_copy_clone_slice_sum_transpose_5(%in_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr1: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr1"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %xnumel: i32 loc("xnumel"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<0> : tensor<32x16xi64, #blocked> loc(#loc1)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc1)
+    %c16_i32 = arith.constant 16 : i32 loc(#loc1)
+    %c1_i64 = arith.constant 1 : i64 loc(#loc1)
+    %c32_i32 = arith.constant 32 : i32 loc(#loc1)
+    %cst_0 = arith.constant dense<0> : tensor<32x16xi32, #blocked> loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc46)
+    %xoffset_1 = arith.muli %xoffset, %c32_i32 : i32 loc(#loc47)
+    %xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc48)
+    %xindex_2 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked> loc(#loc48)
+    %xindex_3 = tt.splat %xoffset_1 : i32 -> tensor<32x1xi32, #blocked> loc(#loc49)
+    %xindex_4 = arith.addi %xindex_3, %xindex_2 : tensor<32x1xi32, #blocked> loc(#loc49)
+    %xmask = tt.splat %xnumel : i32 -> tensor<32x1xi32, #blocked> loc(#loc50)
+    %xmask_5 = arith.cmpi slt, %xindex_4, %xmask : tensor<32x1xi32, #blocked> loc(#loc50)
+    %r0_base = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc51)
+    %r0_base_6 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc51)
+    %x0 = arith.extsi %xindex_4 : tensor<32x1xi32, #blocked> to tensor<32x1xi64, #blocked> loc(#loc52)
+    %x0_7 = tt.splat %ks0 : i64 -> tensor<32x1xi64, #blocked> loc(#loc52)
+    %x0_8 = arith.remsi %x0, %x0_7 : tensor<32x1xi64, #blocked> loc(#loc52)
+    %x1 = arith.divsi %x0, %x0_7 : tensor<32x1xi64, #blocked> loc(#loc53)
+    %r0_mask = tt.splat %r0_numel : i32 -> tensor<1x16xi32, #blocked> loc(#loc54)
+    %tmp0 = tt.splat %ks0 : i64 -> tensor<1x16xi64, #blocked> loc(#loc55)
+    %tmp0_9 = tt.broadcast %x0_8 : tensor<32x1xi64, #blocked> -> tensor<32x16xi64, #blocked> loc(#loc56)
+    %tmp0_10 = arith.muli %ks0, %ks1 : i64 loc(#loc57)
+    %tmp0_11 = tt.splat %tmp0_10 : i64 -> tensor<32x1xi64, #blocked> loc(#loc58)
+    %tmp0_12 = arith.muli %tmp0_11, %x1 : tensor<32x1xi64, #blocked> loc(#loc58)
+    %tmp0_13 = tt.broadcast %tmp0_12 : tensor<32x1xi64, #blocked> -> tensor<32x16xi64, #blocked> loc(#loc59)
+    %tmp0_14 = tt.splat %in_ptr0 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>, #blocked> loc(#loc60)
+    %tmp0_15 = tt.broadcast %xmask_5 : tensor<32x1xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc61)
+    %_tmp3 = scf.for %_tmp3_17 = %c0_i32 to %r0_numel step %c16_i32 iter_args(%_tmp3_18 = %cst) -> (tensor<32x16xi64, #blocked>)  : i32 {
+      %r0_index = tt.splat %_tmp3_17 : i32 -> tensor<1x16xi32, #blocked> loc(#loc63)
+      %r0_index_19 = arith.addi %r0_index, %r0_base_6 : tensor<1x16xi32, #blocked> loc(#loc63)
+      %r0_mask_20 = arith.cmpi slt, %r0_index_19, %r0_mask : tensor<1x16xi32, #blocked> loc(#loc54)
+      %tmp0_21 = arith.extsi %r0_index_19 : tensor<1x16xi32, #blocked> to tensor<1x16xi64, #blocked> loc(#loc55)
+      %tmp0_22 = arith.muli %tmp0, %tmp0_21 : tensor<1x16xi64, #blocked> loc(#loc55)
+      %tmp0_23 = tt.broadcast %tmp0_22 : tensor<1x16xi64, #blocked> -> tensor<32x16xi64, #blocked> loc(#loc56)
+      %tmp0_24 = arith.addi %tmp0_9, %tmp0_23 : tensor<32x16xi64, #blocked> loc(#loc56)
+      %tmp0_25 = arith.addi %tmp0_24, %tmp0_13 : tensor<32x16xi64, #blocked> loc(#loc59)
+      %tmp0_26 = tt.addptr %tmp0_14, %tmp0_25 : tensor<32x16x!tt.ptr<i32>, #blocked>, tensor<32x16xi64, #blocked> loc(#loc60)
+      %tmp0_27 = tt.broadcast %r0_mask_20 : tensor<1x16xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc61)
+      %tmp0_28 = arith.andi %tmp0_27, %tmp0_15 : tensor<32x16xi1, #blocked> loc(#loc61)
+      %tmp0_29 = tt.load %tmp0_26, %tmp0_28, %cst_0 evictionPolicy = evict_last : tensor<32x16x!tt.ptr<i32>, #blocked> loc(#loc64)
+      %tmp1 = arith.extsi %tmp0_29 : tensor<32x16xi32, #blocked> to tensor<32x16xi64, #blocked> loc(#loc65)
+      %tmp4 = arith.addi %_tmp3_18, %tmp1 : tensor<32x16xi64, #blocked> loc(#loc66)
+      %_tmp3_30 = arith.select %tmp0_28, %tmp4, %_tmp3_18 : tensor<32x16xi1, #blocked>, tensor<32x16xi64, #blocked> loc(#loc67)
+      scf.yield %_tmp3_30 : tensor<32x16xi64, #blocked> loc(#loc24)
+    } loc(#loc62)
+    %tmp3 = "tt.reduce"(%_tmp3) <{axis = 1 : i32}> ({
+    ^bb0(%tmp3_17: i64 loc(callsite(#loc1 at #loc68)), %tmp3_18: i64 loc(callsite(#loc1 at #loc68))):
+      %tmp3_19 = arith.addi %tmp3_17, %tmp3_18 : i64 loc(#loc74)
+      tt.reduce.return %tmp3_19 : i64 loc(#loc72)
+    }) : (tensor<32x16xi64, #blocked>) -> tensor<32xi64, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc72)
+    %tmp3_16 = tt.expand_dims %tmp3 {axis = 1 : i32} : tensor<32xi64, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi64, #blocked> loc(#loc69)
+    %tmp5 = arith.trunci %tmp3_16 : tensor<32x1xi64, #blocked> to tensor<32x1xi32, #blocked> loc(#loc70)
+    %0 = arith.cmpi sle, %ks0, %c1_i64 : i64 loc(#loc30)
+    %1 = arith.cmpi sgt, %ks0, %c1_i64 : i64 loc(#loc31)
+    %2 = arith.extui %1 : i1 to i64 loc(#loc32)
+    %3 = arith.muli %ks0, %2 : i64 loc(#loc32)
+    %4 = arith.extui %0 : i1 to i64 loc(#loc71)
+    %5 = arith.addi %4, %3 : i64 loc(#loc33)
+    %6 = tt.splat %5 : i64 -> tensor<32x1xi64, #blocked> loc(#loc35)
+    %7 = arith.muli %x1, %6 : tensor<32x1xi64, #blocked> loc(#loc35)
+    %8 = arith.addi %x0_8, %7 : tensor<32x1xi64, #blocked> loc(#loc36)
+    %9 = tt.splat %out_ptr1 : !tt.ptr<i32> -> tensor<32x1x!tt.ptr<i32>, #blocked> loc(#loc37)
+    %10 = tt.addptr %9, %8 : tensor<32x1x!tt.ptr<i32>, #blocked>, tensor<32x1xi64, #blocked> loc(#loc37)
+    tt.store %10, %tmp5, %xmask_5 : tensor<32x1x!tt.ptr<i32>, #blocked> loc(#loc38)
+    tt.return loc(#loc39)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":21:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":21:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":22:44)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":22:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":23:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":24:37)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":26:19)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":27:19)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":32:29)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:43)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:39)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:54)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:58)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:50)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:34)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:73)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":30:40)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":31:31)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:63)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":37:23)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":39:23)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":40:48)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":40:8)
+#loc25 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc27 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":41:28)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":42:19)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:49)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:75)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:66)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:57)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:41)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:34)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:30)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:25)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:88)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:4)
+#loc46 = loc("xoffset"(#loc2))
+#loc47 = loc("xoffset"(#loc3))
+#loc48 = loc("xindex"(#loc4))
+#loc49 = loc("xindex"(#loc5))
+#loc50 = loc("xmask"(#loc6))
+#loc51 = loc("r0_base"(#loc7))
+#loc52 = loc("x0"(#loc8))
+#loc53 = loc("x1"(#loc9))
+#loc54 = loc("r0_mask"(#loc10))
+#loc55 = loc("tmp0"(#loc11))
+#loc56 = loc("tmp0"(#loc12))
+#loc57 = loc("tmp0"(#loc13))
+#loc58 = loc("tmp0"(#loc14))
+#loc59 = loc("tmp0"(#loc15))
+#loc60 = loc("tmp0"(#loc16))
+#loc61 = loc("tmp0"(#loc17))
+#loc62 = loc("_tmp3"(#loc18))
+#loc63 = loc("r0_index"(#loc19))
+#loc64 = loc("tmp0"(#loc20))
+#loc65 = loc("tmp1"(#loc21))
+#loc66 = loc("tmp4"(#loc22))
+#loc67 = loc("_tmp3"(#loc23))
+#loc69 = loc("tmp3"(#loc28))
+#loc70 = loc("tmp5"(#loc29))
+#loc71 = loc(fused[#loc33, #loc34])
+#loc72 = loc(callsite(#loc25 at #loc68))
+#loc74 = loc(callsite(#loc27 at #loc72))

SpecForge-ext/cache/compiled_kernels/triton/6/2HBOMUT44J5WFCUWYGRFAAS3HGVNDHLHT7HCSXUCAOIKU6XGJNTA/triton_red_fused__to_copy_clone_slice_sum_transpose_5.ttir

@@ -0,0 +1,152 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":18:0)
+#loc1 = loc(unknown)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":41:25)
+#loc43 = loc("in_ptr0"(#loc))
+#loc44 = loc("out_ptr1"(#loc))
+#loc45 = loc("ks0"(#loc))
+#loc46 = loc("ks1"(#loc))
+#loc47 = loc("xnumel"(#loc))
+#loc48 = loc("r0_numel"(#loc))
+#loc74 = loc("tmp3"(#loc29))
+#loc79 = loc(callsite(#loc1 at #loc74))
+module {
+  tt.func public @triton_red_fused__to_copy_clone_slice_sum_transpose_5(%in_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr1: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr1"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %xnumel: i32 loc("xnumel"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %c1_i64 = arith.constant 1 : i64 loc(#loc1)
+    %cst = arith.constant dense<0> : tensor<32x16xi32> loc(#loc1)
+    %c16_i32 = arith.constant 16 : i32 loc(#loc2)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc2)
+    %_tmp3 = arith.constant dense<0> : tensor<32x16xi64> loc(#loc49)
+    %c32_i32 = arith.constant 32 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc50)
+    %xoffset_0 = arith.muli %xoffset, %c32_i32 : i32 loc(#loc51)
+    %xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc52)
+    %xindex_1 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc53)
+    %xindex_2 = tt.splat %xoffset_0 : i32 -> tensor<32x1xi32> loc(#loc54)
+    %xindex_3 = arith.addi %xindex_2, %xindex_1 : tensor<32x1xi32> loc(#loc54)
+    %xmask = tt.splat %xnumel : i32 -> tensor<32x1xi32> loc(#loc55)
+    %xmask_4 = arith.cmpi slt, %xindex_3, %xmask : tensor<32x1xi32> loc(#loc55)
+    %r0_base = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc56)
+    %r0_base_5 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc57)
+    %x0 = arith.extsi %xindex_3 : tensor<32x1xi32> to tensor<32x1xi64> loc(#loc58)
+    %x0_6 = tt.splat %ks0 : i64 -> tensor<32x1xi64> loc(#loc58)
+    %x0_7 = arith.remsi %x0, %x0_6 : tensor<32x1xi64> loc(#loc58)
+    %x1 = arith.divsi %x0, %x0_6 : tensor<32x1xi64> loc(#loc59)
+    %_tmp3_8 = scf.for %r0_offset = %c0_i32 to %r0_numel step %c16_i32 iter_args(%_tmp3_10 = %_tmp3) -> (tensor<32x16xi64>)  : i32 {
+      %r0_index = tt.splat %r0_offset : i32 -> tensor<1x16xi32> loc(#loc61)
+      %r0_index_11 = arith.addi %r0_index, %r0_base_5 : tensor<1x16xi32> loc(#loc61)
+      %r0_mask = tt.splat %r0_numel : i32 -> tensor<1x16xi32> loc(#loc62)
+      %r0_mask_12 = arith.cmpi slt, %r0_index_11, %r0_mask : tensor<1x16xi32> loc(#loc62)
+      %tmp0 = arith.extsi %r0_index_11 : tensor<1x16xi32> to tensor<1x16xi64> loc(#loc63)
+      %tmp0_13 = tt.splat %ks0 : i64 -> tensor<1x16xi64> loc(#loc63)
+      %tmp0_14 = arith.muli %tmp0_13, %tmp0 : tensor<1x16xi64> loc(#loc63)
+      %tmp0_15 = tt.broadcast %x0_7 : tensor<32x1xi64> -> tensor<32x16xi64> loc(#loc64)
+      %tmp0_16 = tt.broadcast %tmp0_14 : tensor<1x16xi64> -> tensor<32x16xi64> loc(#loc64)
+      %tmp0_17 = arith.addi %tmp0_15, %tmp0_16 : tensor<32x16xi64> loc(#loc64)
+      %tmp0_18 = arith.muli %ks0, %ks1 : i64 loc(#loc65)
+      %tmp0_19 = tt.splat %tmp0_18 : i64 -> tensor<32x1xi64> loc(#loc66)
+      %tmp0_20 = arith.muli %tmp0_19, %x1 : tensor<32x1xi64> loc(#loc66)
+      %tmp0_21 = tt.broadcast %tmp0_20 : tensor<32x1xi64> -> tensor<32x16xi64> loc(#loc67)
+      %tmp0_22 = arith.addi %tmp0_17, %tmp0_21 : tensor<32x16xi64> loc(#loc67)
+      %tmp0_23 = tt.splat %in_ptr0 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>> loc(#loc68)
+      %tmp0_24 = tt.addptr %tmp0_23, %tmp0_22 : tensor<32x16x!tt.ptr<i32>>, tensor<32x16xi64> loc(#loc68)
+      %tmp0_25 = tt.broadcast %r0_mask_12 : tensor<1x16xi1> -> tensor<32x16xi1> loc(#loc69)
+      %tmp0_26 = tt.broadcast %xmask_4 : tensor<32x1xi1> -> tensor<32x16xi1> loc(#loc69)
+      %tmp0_27 = arith.andi %tmp0_25, %tmp0_26 : tensor<32x16xi1> loc(#loc69)
+      %tmp0_28 = tt.load %tmp0_24, %tmp0_27, %cst evictionPolicy = evict_last : tensor<32x16x!tt.ptr<i32>> loc(#loc70)
+      %tmp1 = arith.extsi %tmp0_28 : tensor<32x16xi32> to tensor<32x16xi64> loc(#loc71)
+      %tmp4 = arith.addi %_tmp3_10, %tmp1 : tensor<32x16xi64> loc(#loc72)
+      %_tmp3_29 = arith.select %tmp0_27, %tmp4, %_tmp3_10 : tensor<32x16xi1>, tensor<32x16xi64> loc(#loc73)
+      scf.yield %_tmp3_29 : tensor<32x16xi64> loc(#loc27)
+    } loc(#loc60)
+    %tmp3 = "tt.reduce"(%_tmp3_8) <{axis = 1 : i32}> ({
+    ^bb0(%tmp3_10: i64 loc(callsite(#loc1 at #loc74)), %tmp3_11: i64 loc(callsite(#loc1 at #loc74))):
+      %tmp3_12 = arith.addi %tmp3_10, %tmp3_11 : i64 loc(#loc80)
+      tt.reduce.return %tmp3_12 : i64 loc(#loc78)
+    }) : (tensor<32x16xi64>) -> tensor<32xi64> loc(#loc78)
+    %tmp3_9 = tt.expand_dims %tmp3 {axis = 1 : i32} : tensor<32xi64> -> tensor<32x1xi64> loc(#loc75)
+    %tmp5 = arith.trunci %tmp3_9 : tensor<32x1xi64> to tensor<32x1xi32> loc(#loc76)
+    %0 = arith.cmpi sle, %ks0, %c1_i64 : i64 loc(#loc33)
+    %1 = arith.cmpi sgt, %ks0, %c1_i64 : i64 loc(#loc34)
+    %2 = arith.extui %1 : i1 to i64 loc(#loc35)
+    %3 = arith.muli %ks0, %2 : i64 loc(#loc35)
+    %4 = arith.extui %0 : i1 to i64 loc(#loc77)
+    %5 = arith.addi %4, %3 : i64 loc(#loc36)
+    %6 = tt.splat %5 : i64 -> tensor<32x1xi64> loc(#loc38)
+    %7 = arith.muli %x1, %6 : tensor<32x1xi64> loc(#loc38)
+    %8 = arith.addi %x0_7, %7 : tensor<32x1xi64> loc(#loc39)
+    %9 = tt.splat %out_ptr1 : !tt.ptr<i32> -> tensor<32x1x!tt.ptr<i32>> loc(#loc40)
+    %10 = tt.addptr %9, %8 : tensor<32x1x!tt.ptr<i32>>, tensor<32x1xi64> loc(#loc40)
+    tt.store %10, %tmp5, %xmask_4 : tensor<32x1x!tt.ptr<i32>> loc(#loc41)
+    tt.return loc(#loc42)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":30:40)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":28:43)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":21:28)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":21:33)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":22:36)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":22:44)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":22:23)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":23:21)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":24:27)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":24:37)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":26:19)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":27:19)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":31:31)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":32:29)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:43)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:39)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:54)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:58)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:50)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:34)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:73)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":36:63)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":37:23)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":39:23)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":40:48)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":40:8)
+#loc28 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":41:28)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":42:19)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:49)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:75)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:66)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:57)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:41)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:34)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:30)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:25)
+#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:88)
+#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5u/c5ucf2it6fwhmru32vpy2xmwkk4vg7ppnoxmhnkf6ifpb2h2tjsl.py":43:4)
+#loc49 = loc("_tmp3"(#loc3))
+#loc50 = loc("xoffset"(#loc4))
+#loc51 = loc("xoffset"(#loc5))
+#loc52 = loc("xindex"(#loc6))
+#loc53 = loc("xindex"(#loc7))
+#loc54 = loc("xindex"(#loc8))
+#loc55 = loc("xmask"(#loc9))
+#loc56 = loc("r0_base"(#loc10))
+#loc57 = loc("r0_base"(#loc11))
+#loc58 = loc("x0"(#loc12))
+#loc59 = loc("x1"(#loc13))
+#loc60 = loc("_tmp3"(#loc2))
+#loc61 = loc("r0_index"(#loc14))
+#loc62 = loc("r0_mask"(#loc15))
+#loc63 = loc("tmp0"(#loc16))
+#loc64 = loc("tmp0"(#loc17))
+#loc65 = loc("tmp0"(#loc18))
+#loc66 = loc("tmp0"(#loc19))
+#loc67 = loc("tmp0"(#loc20))
+#loc68 = loc("tmp0"(#loc21))
+#loc69 = loc("tmp0"(#loc22))
+#loc70 = loc("tmp0"(#loc23))
+#loc71 = loc("tmp1"(#loc24))
+#loc72 = loc("tmp4"(#loc25))
+#loc73 = loc("_tmp3"(#loc26))
+#loc75 = loc("tmp3"(#loc31))
+#loc76 = loc("tmp5"(#loc32))
+#loc77 = loc(fused[#loc36, #loc37])
+#loc78 = loc(callsite(#loc28 at #loc74))
+#loc80 = loc(callsite(#loc30 at #loc78))

SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/__grp__triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.source", "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttir", "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttgir", "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.llir", "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ptx", "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.cubin", "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json"}}

SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json

@@ -0,0 +1 @@
+{"hash": "e50c6a987142dbb9ad0869ff13cb48e2cad8442b31ebedd676f284f006fc5f0a", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 4, "num_ctas": 1, "num_stages": 1, "warp_size": 32, "maxnreg": null, "cluster_dims": [1, 1, 1], "ptx_version": null, "ptx_options": null, "ir_override": null, "enable_fp_fusion": true, "launch_cooperative_grid": false, "launch_pdl": false, "supported_fp8_dtypes": ["fp8e4b15", "fp8e4nv", "fp8e5"], "deprecated_fp8_dot_operand_dtypes": ["fp8e4b15"], "default_dot_input_precision": "tf32", "allowed_dot_input_precisions": ["tf32", "tf32x3", "ieee"], "max_num_imprecise_acc_default": 1073741824, "extern_libs": [["libdevice", "/workspace/specforge/lib/python3.11/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]], "debug": true, "backend_name": "cuda", "sanitize_overflow": false, "arch": "sm90", "instrumentation_mode": "", "triton_version": "3.5.1", "tensordesc_meta": [], "shared": 2048, "tmem_size": 0, "global_scratch_size": 0, "global_scratch_align": 1, "profile_scratch_size": 0, "profile_scratch_align": 1, "name": "triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3"}

SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttgir

@@ -0,0 +1,841 @@
+#blocked = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [1, 4], order = [0, 1]}>
+#blocked1 = #ttg.blocked<{sizePerThread = [1, 4], threadsPerWarp = [8, 4], warpsPerCTA = [4, 1], order = [1, 0]}>
+#linear = #ttg.linear<{register = [[0, 4], [0, 8]], lane = [[1, 0], [2, 0], [4, 0], [8, 0], [16, 0]], warp = [[0, 1], [0, 2]], block = []}>
+#linear1 = #ttg.linear<{register = [[2, 0, 0], [4, 0, 0]], lane = [[8, 0, 0], [16, 0, 0], [32, 0, 0], [64, 0, 0], [128, 0, 0]], warp = [[0, 1, 0], [1, 0, 0]], block = []}>
+#linear2 = #ttg.linear<{register = [[1, 0, 0], [2, 0, 0]], lane = [[4, 0, 0], [8, 0, 0], [16, 0, 0], [32, 0, 0], [64, 0, 0]], warp = [[0, 0, 1], [0, 1, 0]], block = []}>
+#linear3 = #ttg.linear<{register = [[0, 1, 0], [1, 0, 0]], lane = [[2, 0, 0], [4, 0, 0], [8, 0, 0], [16, 0, 0], [32, 0, 0]], warp = [[0, 0, 1], [0, 0, 2]], block = []}>
+#linear4 = #ttg.linear<{register = [[0, 0, 4], [0, 1, 0]], lane = [[1, 0, 0], [2, 0, 0], [4, 0, 0], [8, 0, 0], [16, 0, 0]], warp = [[0, 0, 1], [0, 0, 2]], block = []}>
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":18:0)
+#loc1 = loc(unknown)
+#loc19 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":662:12)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":41:67)
+#loc24 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":634:73)
+#loc28 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:51)
+#loc33 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:53)
+#loc42 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:50)
+#loc47 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:51)
+#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":45:26)
+#loc77 = loc("in_ptr0"(#loc))
+#loc78 = loc("out_ptr2"(#loc))
+#loc79 = loc("out_ptr3"(#loc))
+#loc80 = loc("xnumel"(#loc))
+#loc81 = loc("r0_numel"(#loc))
+#loc99 = loc(callsite(#loc19 at #loc20))
+#loc105 = loc("ileft"(#loc28))
+#loc109 = loc("iright"(#loc33))
+#loc118 = loc("left_idx"(#loc42))
+#loc123 = loc("right_idx"(#loc47))
+#loc143 = loc("tmp11"(#loc67))
+#loc149 = loc(callsite(#loc24 at #loc99))
+#loc153 = loc(callsite(#loc1 at #loc143))
+#loc157 = loc(callsite(#loc105 at #loc149))
+#loc161 = loc(callsite(#loc109 at #loc149))
+#loc169 = loc(callsite(#loc118 at #loc149))
+#loc174 = loc(callsite(#loc123 at #loc149))
+#loc194 = loc(callsite(#loc1 at #loc157))
+#loc196 = loc(callsite(#loc1 at #loc161))
+#loc199 = loc(callsite(#loc1 at #loc169))
+#loc202 = loc(callsite(#loc1 at #loc174))
+module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
+  tt.func public @triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3(%in_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr2: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr2"(#loc)), %out_ptr3: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr3"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<0> : tensor<32x16xi32, #linear> loc(#loc1)
+    %cst_0 = arith.constant dense<0> : tensor<32x16xi64, #blocked> loc(#loc1)
+    %c32_i32 = arith.constant 32 : i32 loc(#loc1)
+    %cst_1 = arith.constant dense<128> : tensor<32x1xi32, #blocked> loc(#loc1)
+    %cst_2 = arith.constant dense<128> : tensor<32x1xi32, #blocked1> loc(#loc1)
+    %cst_3 = arith.constant dense<16> : tensor<32x1xi32, #blocked> loc(#loc1)
+    %cst_4 = arith.constant dense<16> : tensor<32x1xi32, #blocked1> loc(#loc1)
+    %cst_5 = arith.constant dense<17> : tensor<1x16xi32, #blocked> loc(#loc1)
+    %cst_6 = arith.constant dense<272> : tensor<32x1xi32, #blocked> loc(#loc1)
+    %cst_7 = arith.constant dense<1> : tensor<1x2x1xi32, #linear1> loc(#loc1)
+    %cst_8 = arith.constant dense<1> : tensor<1x2x1xi32, #linear2> loc(#loc1)
+    %cst_9 = arith.constant dense<1> : tensor<1x2x1xi32, #linear3> loc(#loc1)
+    %cst_10 = arith.constant dense<1> : tensor<1x2x1xi32, #linear4> loc(#loc1)
+    %cst_11 = arith.constant dense<0> : tensor<32x16xi32, #blocked> loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc82)
+    %xoffset_12 = arith.muli %xoffset, %c32_i32 : i32 loc(#loc83)
+    %xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc84)
+    %xindex_13 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> loc(#loc84)
+    %xindex_14 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked> loc(#loc84)
+    %xindex_15 = tt.expand_dims %xindex_13 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> -> tensor<32x1xi32, #blocked1> loc(#loc84)
+    %xindex_16 = tt.splat %xoffset_12 : i32 -> tensor<32x1xi32, #blocked> loc(#loc85)
+    %xindex_17 = tt.splat %xoffset_12 : i32 -> tensor<32x1xi32, #blocked1> loc(#loc85)
+    %xindex_18 = arith.addi %xindex_16, %xindex_14 : tensor<32x1xi32, #blocked> loc(#loc85)
+    %xindex_19 = arith.addi %xindex_17, %xindex_15 : tensor<32x1xi32, #blocked1> loc(#loc85)
+    %xmask = arith.cmpi slt, %xindex_18, %cst_1 : tensor<32x1xi32, #blocked> loc(#loc86)
+    %xmask_20 = arith.cmpi slt, %xindex_19, %cst_2 : tensor<32x1xi32, #blocked1> loc(#loc86)
+    %r0_index = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc87)
+    %r0_index_21 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #linear}>> loc(#loc87)
+    %r0_index_22 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked1}>> loc(#loc87)
+    %r0_index_23 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc87)
+    %r0_index_24 = tt.expand_dims %r0_index_21 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #linear}>> -> tensor<1x16xi32, #linear> loc(#loc87)
+    %r0_index_25 = tt.expand_dims %r0_index_22 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked1}>> -> tensor<1x16xi32, #blocked1> loc(#loc87)
+    %x0 = arith.remsi %xindex_18, %cst_3 : tensor<32x1xi32, #blocked> loc(#loc88)
+    %x1 = arith.divsi %xindex_18, %cst_3 : tensor<32x1xi32, #blocked> loc(#loc89)
+    %tmp0 = arith.muli %r0_index_23, %cst_5 : tensor<1x16xi32, #blocked> loc(#loc90)
+    %tmp0_26 = tt.broadcast %x0 : tensor<32x1xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc91)
+    %tmp0_27 = tt.broadcast %tmp0 : tensor<1x16xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc91)
+    %tmp0_28 = arith.addi %tmp0_26, %tmp0_27 : tensor<32x16xi32, #blocked> loc(#loc91)
+    %tmp0_29 = arith.muli %x1, %cst_6 : tensor<32x1xi32, #blocked> loc(#loc92)
+    %tmp0_30 = tt.broadcast %tmp0_29 : tensor<32x1xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc93)
+    %tmp0_31 = arith.addi %tmp0_28, %tmp0_30 : tensor<32x16xi32, #blocked> loc(#loc93)
+    %tmp0_32 = tt.splat %in_ptr0 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>, #blocked> loc(#loc94)
+    %tmp0_33 = tt.addptr %tmp0_32, %tmp0_31 : tensor<32x16x!tt.ptr<i32>, #blocked>, tensor<32x16xi32, #blocked> loc(#loc94)
+    %tmp0_34 = tt.broadcast %xmask : tensor<32x1xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc95)
+    %tmp0_35 = tt.broadcast %xmask_20 : tensor<32x1xi1, #blocked1> -> tensor<32x16xi1, #blocked1> loc(#loc95)
+    %tmp0_36 = tt.load %tmp0_33, %tmp0_34, %cst_11 : tensor<32x16x!tt.ptr<i32>, #blocked> loc(#loc95)
+    %tmp2 = arith.trunci %r0_index_24 : tensor<1x16xi32, #linear> to tensor<1x16xi16, #linear> loc(#loc96)
+    %tmp4 = tt.broadcast %tmp2 : tensor<1x16xi16, #linear> -> tensor<32x16xi16, #linear> loc(#loc97)
+    %flip = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear2}>}>> loc(#loc146)
+    %flip_37 = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear1}>}>> loc(#loc146)
+    %flip_38 = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear3}>}>> loc(#loc146)
+    %flip_39 = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear4}>}>> loc(#loc146)
+    %flip_40 = tt.expand_dims %flip {axis = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear2}>}>> -> tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear2}>> loc(#loc146)
+    %flip_41 = tt.expand_dims %flip_37 {axis = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear1}>}>> -> tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear1}>> loc(#loc146)
+    %flip_42 = tt.expand_dims %flip_38 {axis = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear3}>}>> -> tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear3}>> loc(#loc146)
+    %flip_43 = tt.expand_dims %flip_39 {axis = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #ttg.slice<{dim = 2, parent = #linear4}>}>> -> tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear4}>> loc(#loc146)
+    %flip_44 = tt.expand_dims %flip_40 {axis = 2 : i32} : tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear2}>> -> tensor<1x2x1xi32, #linear2> loc(#loc146)
+    %flip_45 = tt.expand_dims %flip_41 {axis = 2 : i32} : tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear1}>> -> tensor<1x2x1xi32, #linear1> loc(#loc146)
+    %flip_46 = tt.expand_dims %flip_42 {axis = 2 : i32} : tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear3}>> -> tensor<1x2x1xi32, #linear3> loc(#loc146)
+    %flip_47 = tt.expand_dims %flip_43 {axis = 2 : i32} : tensor<1x2xi32, #ttg.slice<{dim = 2, parent = #linear4}>> -> tensor<1x2x1xi32, #linear4> loc(#loc146)
+    %flip_48 = tt.broadcast %flip_44 : tensor<1x2x1xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc147)
+    %flip_49 = tt.reshape %flip_48 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #blocked> loc(#loc148)
+    %flip_50 = tt.reshape %flip_48 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc148)
+    %y = tt.reshape %tmp0_36 : tensor<32x16xi32, #blocked> -> tensor<256x2x1xi32, #linear1> loc(#loc154)
+    %left_mask = arith.subi %cst_7, %flip_45 : tensor<1x2x1xi32, #linear1> loc(#loc155)
+    %left_mask_51 = arith.subi %cst_8, %flip_44 : tensor<1x2x1xi32, #linear2> loc(#loc155)
+    %left_mask_52 = arith.subi %cst_9, %flip_46 : tensor<1x2x1xi32, #linear3> loc(#loc155)
+    %left_mask_53 = arith.subi %cst_10, %flip_47 : tensor<1x2x1xi32, #linear4> loc(#loc155)
+    %ileft = tt.broadcast %left_mask : tensor<1x2x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc156)
+    %ileft_54 = arith.muli %y, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc156)
+    %ileft_55 = "tt.reduce"(%ileft_54) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc193)
+    %ileft_56 = tt.expand_dims %ileft_55 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc158)
+    %ileft_57 = tt.broadcast %ileft_56 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc159)
+    %iright = tt.broadcast %flip_45 : tensor<1x2x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc160)
+    %iright_58 = arith.muli %y, %iright : tensor<256x2x1xi32, #linear1> loc(#loc160)
+    %iright_59 = "tt.reduce"(%iright_58) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc195)
+    %iright_60 = tt.expand_dims %iright_59 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc162)
+    %iright_61 = tt.broadcast %iright_60 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc163)
+    %ileft_62 = tt.reshape %ileft_57 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #blocked> loc(#loc164)
+    %ileft_63 = tt.reshape %ileft_57 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_64 = tt.reshape %iright_61 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #blocked> loc(#loc165)
+    %iright_65 = tt.reshape %iright_61 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx = tt.reshape %tmp4 : tensor<32x16xi16, #linear> -> tensor<256x2x1xi16, #linear1> loc(#loc166)
+    %left_idx = arith.trunci %left_mask : tensor<1x2x1xi32, #linear1> to tensor<1x2x1xi16, #linear1> loc(#loc167)
+    %left_idx_66 = tt.broadcast %left_idx : tensor<1x2x1xi16, #linear1> -> tensor<256x2x1xi16, #linear1> loc(#loc168)
+    %left_idx_67 = arith.muli %y_idx, %left_idx_66 : tensor<256x2x1xi16, #linear1> loc(#loc168)
+    %input = arith.extsi %left_idx_67 : tensor<256x2x1xi16, #linear1> to tensor<256x2x1xi32, #linear1> loc(#loc197)
+    %left_idx_68 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc198)
+    %left_idx_69 = tt.expand_dims %left_idx_68 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc170)
+    %left_idx_70 = tt.broadcast %left_idx_69 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc171)
+    %right_idx = arith.trunci %flip_45 : tensor<1x2x1xi32, #linear1> to tensor<1x2x1xi16, #linear1> loc(#loc172)
+    %right_idx_71 = tt.broadcast %right_idx : tensor<1x2x1xi16, #linear1> -> tensor<256x2x1xi16, #linear1> loc(#loc173)
+    %right_idx_72 = arith.muli %y_idx, %right_idx_71 : tensor<256x2x1xi16, #linear1> loc(#loc173)
+    %input_73 = arith.extsi %right_idx_72 : tensor<256x2x1xi16, #linear1> to tensor<256x2x1xi32, #linear1> loc(#loc200)
+    %right_idx_74 = "tt.reduce"(%input_73) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc201)
+    %right_idx_75 = tt.expand_dims %right_idx_74 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc175)
+    %right_idx_76 = tt.broadcast %right_idx_75 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc176)
+    %left_idx_77 = tt.reshape %left_idx_70 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #blocked> loc(#loc177)
+    %left_idx_78 = tt.reshape %left_idx_70 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_79 = tt.reshape %right_idx_76 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #blocked> loc(#loc178)
+    %right_idx_80 = tt.reshape %right_idx_76 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond = arith.cmpi slt, %ileft_62, %iright_64 : tensor<32x16xi32, #blocked> loc(#loc179)
+    %cond_81 = arith.cmpi slt, %ileft_63, %iright_65 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq = arith.cmpi eq, %ileft_62, %iright_64 : tensor<32x16xi32, #blocked> loc(#loc180)
+    %eq_82 = arith.cmpi eq, %ileft_63, %iright_65 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_83 = arith.cmpi sgt, %left_idx_77, %right_idx_79 : tensor<32x16xi32, #blocked> loc(#loc181)
+    %cond_84 = arith.cmpi sgt, %left_idx_78, %right_idx_80 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_85 = arith.andi %eq, %cond_83 : tensor<32x16xi1, #blocked> loc(#loc182)
+    %cond_86 = arith.andi %eq_82, %cond_84 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_87 = arith.ori %cond, %cond_85 : tensor<32x16xi1, #blocked> loc(#loc183)
+    %cond_88 = arith.ori %cond_81, %cond_86 : tensor<32x16xi1, #linear> loc(#loc183)
+    %cond_89 = arith.extui %cond_87 : tensor<32x16xi1, #blocked> to tensor<32x16xi32, #blocked> loc(#loc184)
+    %cond_90 = arith.extui %cond_88 : tensor<32x16xi1, #linear> to tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_91 = arith.xori %cond_89, %flip_49 : tensor<32x16xi32, #blocked> loc(#loc184)
+    %cond_92 = arith.xori %cond_90, %flip_50 : tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_93 = arith.cmpi ne, %cond_91, %cst_11 : tensor<32x16xi32, #blocked> loc(#loc185)
+    %cond_94 = arith.cmpi ne, %cond_92, %cst : tensor<32x16xi32, #linear> loc(#loc185)
+    %ret = arith.xori %ileft_62, %iright_64 : tensor<32x16xi32, #blocked> loc(#loc186)
+    %ret_95 = arith.select %cond_93, %ret, %cst_11 : tensor<32x16xi1, #blocked>, tensor<32x16xi32, #blocked> loc(#loc187)
+    %ret_96 = arith.xori %tmp0_36, %ret_95 : tensor<32x16xi32, #blocked> loc(#loc188)
+    %ret_97 = ttg.convert_layout %ret_96 : tensor<32x16xi32, #blocked> -> tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs = arith.xori %left_idx_78, %right_idx_80 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_98 = arith.select %cond_94, %new_idxs, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_99 = arith.extsi %tmp2 : tensor<1x16xi16, #linear> to tensor<1x16xi32, #linear> loc(#loc191)
+    %new_idxs_100 = tt.broadcast %new_idxs_99 : tensor<1x16xi32, #linear> -> tensor<32x16xi32, #linear> loc(#loc191)
+    %new_idxs_101 = arith.xori %new_idxs_100, %new_idxs_98 : tensor<32x16xi32, #linear> loc(#loc191)
+    %flip_102 = tt.broadcast %flip_46 : tensor<1x2x1xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc147)
+    %flip_103 = tt.reshape %flip_102 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc148)
+    %y_104 = tt.reshape %ret_96 : tensor<32x16xi32, #blocked> -> tensor<128x2x2xi32, #linear2> loc(#loc154)
+    %ileft_105 = tt.broadcast %left_mask_51 : tensor<1x2x1xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc156)
+    %ileft_106 = arith.muli %y_104, %ileft_105 : tensor<128x2x2xi32, #linear2> loc(#loc156)
+    %ileft_107 = "tt.reduce"(%ileft_106) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc193)
+    %ileft_108 = tt.expand_dims %ileft_107 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc158)
+    %ileft_109 = tt.broadcast %ileft_108 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc159)
+    %iright_110 = arith.muli %y_104, %flip_48 : tensor<128x2x2xi32, #linear2> loc(#loc160)
+    %iright_111 = "tt.reduce"(%iright_110) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc195)
+    %iright_112 = tt.expand_dims %iright_111 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc162)
+    %iright_113 = tt.broadcast %iright_112 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc163)
+    %ileft_114 = tt.reshape %ileft_109 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_115 = tt.reshape %iright_113 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_116 = tt.reshape %new_idxs_101 : tensor<32x16xi32, #linear> -> tensor<128x2x2xi32, #linear2> loc(#loc166)
+    %left_idx_117 = arith.muli %y_idx_116, %ileft_105 : tensor<128x2x2xi32, #linear2> loc(#loc168)
+    %left_idx_118 = "tt.reduce"(%left_idx_117) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc198)
+    %left_idx_119 = tt.expand_dims %left_idx_118 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc170)
+    %left_idx_120 = tt.broadcast %left_idx_119 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc171)
+    %right_idx_121 = arith.muli %y_idx_116, %flip_48 : tensor<128x2x2xi32, #linear2> loc(#loc173)
+    %right_idx_122 = "tt.reduce"(%right_idx_121) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc201)
+    %right_idx_123 = tt.expand_dims %right_idx_122 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc175)
+    %right_idx_124 = tt.broadcast %right_idx_123 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc176)
+    %left_idx_125 = tt.reshape %left_idx_120 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_126 = tt.reshape %right_idx_124 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_127 = arith.cmpi slt, %ileft_114, %iright_115 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_128 = arith.cmpi eq, %ileft_114, %iright_115 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_129 = arith.cmpi sgt, %left_idx_125, %right_idx_126 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_130 = arith.andi %eq_128, %cond_129 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_131 = arith.ori %cond_127, %cond_130 : tensor<32x16xi1, #linear> loc(#loc183)
+    %cond_132 = arith.extui %cond_131 : tensor<32x16xi1, #linear> to tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_133 = arith.xori %cond_132, %flip_103 : tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_134 = arith.cmpi ne, %cond_133, %cst : tensor<32x16xi32, #linear> loc(#loc185)
+    %ret_135 = arith.xori %ileft_114, %iright_115 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_136 = arith.select %cond_134, %ret_135, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_137 = arith.xori %ret_97, %ret_136 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_138 = arith.xori %left_idx_125, %right_idx_126 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_139 = arith.select %cond_134, %new_idxs_138, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_140 = arith.xori %new_idxs_101, %new_idxs_139 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_141 = tt.reshape %ret_137 : tensor<32x16xi32, #linear> -> tensor<256x2x1xi32, #linear1> loc(#loc154)
+    %ileft_142 = arith.muli %y_141, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc156)
+    %ileft_143 = "tt.reduce"(%ileft_142) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc193)
+    %ileft_144 = tt.expand_dims %ileft_143 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc158)
+    %ileft_145 = tt.broadcast %ileft_144 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc159)
+    %iright_146 = arith.muli %y_141, %iright : tensor<256x2x1xi32, #linear1> loc(#loc160)
+    %iright_147 = "tt.reduce"(%iright_146) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc195)
+    %iright_148 = tt.expand_dims %iright_147 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc162)
+    %iright_149 = tt.broadcast %iright_148 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc163)
+    %ileft_150 = tt.reshape %ileft_145 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_151 = tt.reshape %iright_149 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_152 = tt.reshape %new_idxs_140 : tensor<32x16xi32, #linear> -> tensor<256x2x1xi32, #linear1> loc(#loc166)
+    %left_idx_153 = arith.muli %y_idx_152, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc168)
+    %left_idx_154 = "tt.reduce"(%left_idx_153) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc198)
+    %left_idx_155 = tt.expand_dims %left_idx_154 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc170)
+    %left_idx_156 = tt.broadcast %left_idx_155 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc171)
+    %right_idx_157 = arith.muli %y_idx_152, %iright : tensor<256x2x1xi32, #linear1> loc(#loc173)
+    %right_idx_158 = "tt.reduce"(%right_idx_157) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc201)
+    %right_idx_159 = tt.expand_dims %right_idx_158 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc175)
+    %right_idx_160 = tt.broadcast %right_idx_159 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc176)
+    %left_idx_161 = tt.reshape %left_idx_156 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_162 = tt.reshape %right_idx_160 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_163 = arith.cmpi slt, %ileft_150, %iright_151 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_164 = arith.cmpi eq, %ileft_150, %iright_151 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_165 = arith.cmpi sgt, %left_idx_161, %right_idx_162 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_166 = arith.andi %eq_164, %cond_165 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_167 = arith.ori %cond_163, %cond_166 : tensor<32x16xi1, #linear> loc(#loc183)
+    %cond_168 = arith.extui %cond_167 : tensor<32x16xi1, #linear> to tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_169 = arith.xori %cond_168, %flip_103 : tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_170 = arith.cmpi ne, %cond_169, %cst : tensor<32x16xi32, #linear> loc(#loc185)
+    %ret_171 = arith.xori %ileft_150, %iright_151 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_172 = arith.select %cond_170, %ret_171, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_173 = arith.xori %ret_137, %ret_172 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_174 = arith.xori %left_idx_161, %right_idx_162 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_175 = arith.select %cond_170, %new_idxs_174, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_176 = arith.xori %new_idxs_140, %new_idxs_175 : tensor<32x16xi32, #linear> loc(#loc191)
+    %flip_177 = tt.broadcast %flip_47 : tensor<1x2x1xi32, #linear4> -> tensor<32x2x8xi32, #linear4> loc(#loc147)
+    %flip_178 = tt.reshape %flip_177 : tensor<32x2x8xi32, #linear4> -> tensor<32x16xi32, #linear> loc(#loc148)
+    %y_179 = tt.reshape %ret_173 : tensor<32x16xi32, #linear> -> tensor<64x2x4xi32, #linear3> loc(#loc154)
+    %ileft_180 = tt.broadcast %left_mask_52 : tensor<1x2x1xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc156)
+    %ileft_181 = arith.muli %y_179, %ileft_180 : tensor<64x2x4xi32, #linear3> loc(#loc156)
+    %ileft_182 = "tt.reduce"(%ileft_181) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc193)
+    %ileft_183 = tt.expand_dims %ileft_182 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc158)
+    %ileft_184 = tt.broadcast %ileft_183 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc159)
+    %iright_185 = arith.muli %y_179, %flip_102 : tensor<64x2x4xi32, #linear3> loc(#loc160)
+    %iright_186 = "tt.reduce"(%iright_185) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc195)
+    %iright_187 = tt.expand_dims %iright_186 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc162)
+    %iright_188 = tt.broadcast %iright_187 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc163)
+    %ileft_189 = tt.reshape %ileft_184 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_190 = tt.reshape %iright_188 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_191 = tt.reshape %new_idxs_176 : tensor<32x16xi32, #linear> -> tensor<64x2x4xi32, #linear3> loc(#loc166)
+    %left_idx_192 = arith.muli %y_idx_191, %ileft_180 : tensor<64x2x4xi32, #linear3> loc(#loc168)
+    %left_idx_193 = "tt.reduce"(%left_idx_192) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc198)
+    %left_idx_194 = tt.expand_dims %left_idx_193 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc170)
+    %left_idx_195 = tt.broadcast %left_idx_194 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc171)
+    %right_idx_196 = arith.muli %y_idx_191, %flip_102 : tensor<64x2x4xi32, #linear3> loc(#loc173)
+    %right_idx_197 = "tt.reduce"(%right_idx_196) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc201)
+    %right_idx_198 = tt.expand_dims %right_idx_197 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc175)
+    %right_idx_199 = tt.broadcast %right_idx_198 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc176)
+    %left_idx_200 = tt.reshape %left_idx_195 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_201 = tt.reshape %right_idx_199 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_202 = arith.cmpi slt, %ileft_189, %iright_190 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_203 = arith.cmpi eq, %ileft_189, %iright_190 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_204 = arith.cmpi sgt, %left_idx_200, %right_idx_201 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_205 = arith.andi %eq_203, %cond_204 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_206 = arith.ori %cond_202, %cond_205 : tensor<32x16xi1, #linear> loc(#loc183)
+    %cond_207 = arith.extui %cond_206 : tensor<32x16xi1, #linear> to tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_208 = arith.xori %cond_207, %flip_178 : tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_209 = arith.cmpi ne, %cond_208, %cst : tensor<32x16xi32, #linear> loc(#loc185)
+    %ret_210 = arith.xori %ileft_189, %iright_190 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_211 = arith.select %cond_209, %ret_210, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_212 = arith.xori %ret_173, %ret_211 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_213 = arith.xori %left_idx_200, %right_idx_201 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_214 = arith.select %cond_209, %new_idxs_213, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_215 = arith.xori %new_idxs_176, %new_idxs_214 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_216 = tt.reshape %ret_212 : tensor<32x16xi32, #linear> -> tensor<128x2x2xi32, #linear2> loc(#loc154)
+    %ileft_217 = arith.muli %y_216, %ileft_105 : tensor<128x2x2xi32, #linear2> loc(#loc156)
+    %ileft_218 = "tt.reduce"(%ileft_217) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc193)
+    %ileft_219 = tt.expand_dims %ileft_218 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc158)
+    %ileft_220 = tt.broadcast %ileft_219 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc159)
+    %iright_221 = arith.muli %y_216, %flip_48 : tensor<128x2x2xi32, #linear2> loc(#loc160)
+    %iright_222 = "tt.reduce"(%iright_221) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc195)
+    %iright_223 = tt.expand_dims %iright_222 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc162)
+    %iright_224 = tt.broadcast %iright_223 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc163)
+    %ileft_225 = tt.reshape %ileft_220 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_226 = tt.reshape %iright_224 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_227 = tt.reshape %new_idxs_215 : tensor<32x16xi32, #linear> -> tensor<128x2x2xi32, #linear2> loc(#loc166)
+    %left_idx_228 = arith.muli %y_idx_227, %ileft_105 : tensor<128x2x2xi32, #linear2> loc(#loc168)
+    %left_idx_229 = "tt.reduce"(%left_idx_228) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc198)
+    %left_idx_230 = tt.expand_dims %left_idx_229 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc170)
+    %left_idx_231 = tt.broadcast %left_idx_230 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc171)
+    %right_idx_232 = arith.muli %y_idx_227, %flip_48 : tensor<128x2x2xi32, #linear2> loc(#loc173)
+    %right_idx_233 = "tt.reduce"(%right_idx_232) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc201)
+    %right_idx_234 = tt.expand_dims %right_idx_233 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc175)
+    %right_idx_235 = tt.broadcast %right_idx_234 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc176)
+    %left_idx_236 = tt.reshape %left_idx_231 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_237 = tt.reshape %right_idx_235 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_238 = arith.cmpi slt, %ileft_225, %iright_226 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_239 = arith.cmpi eq, %ileft_225, %iright_226 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_240 = arith.cmpi sgt, %left_idx_236, %right_idx_237 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_241 = arith.andi %eq_239, %cond_240 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_242 = arith.ori %cond_238, %cond_241 : tensor<32x16xi1, #linear> loc(#loc183)
+    %cond_243 = arith.extui %cond_242 : tensor<32x16xi1, #linear> to tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_244 = arith.xori %cond_243, %flip_178 : tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_245 = arith.cmpi ne, %cond_244, %cst : tensor<32x16xi32, #linear> loc(#loc185)
+    %ret_246 = arith.xori %ileft_225, %iright_226 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_247 = arith.select %cond_245, %ret_246, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_248 = arith.xori %ret_212, %ret_247 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_249 = arith.xori %left_idx_236, %right_idx_237 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_250 = arith.select %cond_245, %new_idxs_249, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_251 = arith.xori %new_idxs_215, %new_idxs_250 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_252 = tt.reshape %ret_248 : tensor<32x16xi32, #linear> -> tensor<256x2x1xi32, #linear1> loc(#loc154)
+    %ileft_253 = arith.muli %y_252, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc156)
+    %ileft_254 = "tt.reduce"(%ileft_253) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc193)
+    %ileft_255 = tt.expand_dims %ileft_254 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc158)
+    %ileft_256 = tt.broadcast %ileft_255 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc159)
+    %iright_257 = arith.muli %y_252, %iright : tensor<256x2x1xi32, #linear1> loc(#loc160)
+    %iright_258 = "tt.reduce"(%iright_257) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc195)
+    %iright_259 = tt.expand_dims %iright_258 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc162)
+    %iright_260 = tt.broadcast %iright_259 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc163)
+    %ileft_261 = tt.reshape %ileft_256 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_262 = tt.reshape %iright_260 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_263 = tt.reshape %new_idxs_251 : tensor<32x16xi32, #linear> -> tensor<256x2x1xi32, #linear1> loc(#loc166)
+    %left_idx_264 = arith.muli %y_idx_263, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc168)
+    %left_idx_265 = "tt.reduce"(%left_idx_264) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc198)
+    %left_idx_266 = tt.expand_dims %left_idx_265 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc170)
+    %left_idx_267 = tt.broadcast %left_idx_266 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc171)
+    %right_idx_268 = arith.muli %y_idx_263, %iright : tensor<256x2x1xi32, #linear1> loc(#loc173)
+    %right_idx_269 = "tt.reduce"(%right_idx_268) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc201)
+    %right_idx_270 = tt.expand_dims %right_idx_269 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc175)
+    %right_idx_271 = tt.broadcast %right_idx_270 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc176)
+    %left_idx_272 = tt.reshape %left_idx_267 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_273 = tt.reshape %right_idx_271 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_274 = arith.cmpi slt, %ileft_261, %iright_262 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_275 = arith.cmpi eq, %ileft_261, %iright_262 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_276 = arith.cmpi sgt, %left_idx_272, %right_idx_273 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_277 = arith.andi %eq_275, %cond_276 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_278 = arith.ori %cond_274, %cond_277 : tensor<32x16xi1, #linear> loc(#loc183)
+    %cond_279 = arith.extui %cond_278 : tensor<32x16xi1, #linear> to tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_280 = arith.xori %cond_279, %flip_178 : tensor<32x16xi32, #linear> loc(#loc184)
+    %cond_281 = arith.cmpi ne, %cond_280, %cst : tensor<32x16xi32, #linear> loc(#loc185)
+    %ret_282 = arith.xori %ileft_261, %iright_262 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_283 = arith.select %cond_281, %ret_282, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_284 = arith.xori %ret_248, %ret_283 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_285 = arith.xori %left_idx_272, %right_idx_273 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_286 = arith.select %cond_281, %new_idxs_285, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_287 = arith.xori %new_idxs_251, %new_idxs_286 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_288 = tt.reshape %ret_284 : tensor<32x16xi32, #linear> -> tensor<32x2x8xi32, #linear4> loc(#loc154)
+    %ileft_289 = tt.broadcast %left_mask_53 : tensor<1x2x1xi32, #linear4> -> tensor<32x2x8xi32, #linear4> loc(#loc156)
+    %ileft_290 = arith.muli %y_288, %ileft_289 : tensor<32x2x8xi32, #linear4> loc(#loc156)
+    %ileft_291 = "tt.reduce"(%ileft_290) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<32x2x8xi32, #linear4>) -> tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> loc(#loc193)
+    %ileft_292 = tt.expand_dims %ileft_291 {axis = 1 : i32} : tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> -> tensor<32x1x8xi32, #linear4> loc(#loc158)
+    %ileft_293 = tt.broadcast %ileft_292 : tensor<32x1x8xi32, #linear4> -> tensor<32x2x8xi32, #linear4> loc(#loc159)
+    %iright_294 = arith.muli %y_288, %flip_177 : tensor<32x2x8xi32, #linear4> loc(#loc160)
+    %iright_295 = "tt.reduce"(%iright_294) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<32x2x8xi32, #linear4>) -> tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> loc(#loc195)
+    %iright_296 = tt.expand_dims %iright_295 {axis = 1 : i32} : tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> -> tensor<32x1x8xi32, #linear4> loc(#loc162)
+    %iright_297 = tt.broadcast %iright_296 : tensor<32x1x8xi32, #linear4> -> tensor<32x2x8xi32, #linear4> loc(#loc163)
+    %ileft_298 = tt.reshape %ileft_293 : tensor<32x2x8xi32, #linear4> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_299 = tt.reshape %iright_297 : tensor<32x2x8xi32, #linear4> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_300 = tt.reshape %new_idxs_287 : tensor<32x16xi32, #linear> -> tensor<32x2x8xi32, #linear4> loc(#loc166)
+    %left_idx_301 = arith.muli %y_idx_300, %ileft_289 : tensor<32x2x8xi32, #linear4> loc(#loc168)
+    %left_idx_302 = "tt.reduce"(%left_idx_301) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<32x2x8xi32, #linear4>) -> tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> loc(#loc198)
+    %left_idx_303 = tt.expand_dims %left_idx_302 {axis = 1 : i32} : tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> -> tensor<32x1x8xi32, #linear4> loc(#loc170)
+    %left_idx_304 = tt.broadcast %left_idx_303 : tensor<32x1x8xi32, #linear4> -> tensor<32x2x8xi32, #linear4> loc(#loc171)
+    %right_idx_305 = arith.muli %y_idx_300, %flip_177 : tensor<32x2x8xi32, #linear4> loc(#loc173)
+    %right_idx_306 = "tt.reduce"(%right_idx_305) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<32x2x8xi32, #linear4>) -> tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> loc(#loc201)
+    %right_idx_307 = tt.expand_dims %right_idx_306 {axis = 1 : i32} : tensor<32x8xi32, #ttg.slice<{dim = 1, parent = #linear4}>> -> tensor<32x1x8xi32, #linear4> loc(#loc175)
+    %right_idx_308 = tt.broadcast %right_idx_307 : tensor<32x1x8xi32, #linear4> -> tensor<32x2x8xi32, #linear4> loc(#loc176)
+    %left_idx_309 = tt.reshape %left_idx_304 : tensor<32x2x8xi32, #linear4> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_310 = tt.reshape %right_idx_308 : tensor<32x2x8xi32, #linear4> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_311 = arith.cmpi slt, %ileft_298, %iright_299 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_312 = arith.cmpi eq, %ileft_298, %iright_299 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_313 = arith.cmpi sgt, %left_idx_309, %right_idx_310 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_314 = arith.andi %eq_312, %cond_313 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_315 = arith.ori %cond_311, %cond_314 : tensor<32x16xi1, #linear> loc(#loc183)
+    %ret_316 = arith.xori %ileft_298, %iright_299 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_317 = arith.select %cond_315, %ret_316, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_318 = arith.xori %ret_284, %ret_317 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_319 = arith.xori %left_idx_309, %right_idx_310 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_320 = arith.select %cond_315, %new_idxs_319, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_321 = arith.xori %new_idxs_287, %new_idxs_320 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_322 = tt.reshape %ret_318 : tensor<32x16xi32, #linear> -> tensor<64x2x4xi32, #linear3> loc(#loc154)
+    %ileft_323 = arith.muli %y_322, %ileft_180 : tensor<64x2x4xi32, #linear3> loc(#loc156)
+    %ileft_324 = "tt.reduce"(%ileft_323) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc193)
+    %ileft_325 = tt.expand_dims %ileft_324 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc158)
+    %ileft_326 = tt.broadcast %ileft_325 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc159)
+    %iright_327 = arith.muli %y_322, %flip_102 : tensor<64x2x4xi32, #linear3> loc(#loc160)
+    %iright_328 = "tt.reduce"(%iright_327) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc195)
+    %iright_329 = tt.expand_dims %iright_328 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc162)
+    %iright_330 = tt.broadcast %iright_329 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc163)
+    %ileft_331 = tt.reshape %ileft_326 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_332 = tt.reshape %iright_330 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_333 = tt.reshape %new_idxs_321 : tensor<32x16xi32, #linear> -> tensor<64x2x4xi32, #linear3> loc(#loc166)
+    %left_idx_334 = arith.muli %y_idx_333, %ileft_180 : tensor<64x2x4xi32, #linear3> loc(#loc168)
+    %left_idx_335 = "tt.reduce"(%left_idx_334) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc198)
+    %left_idx_336 = tt.expand_dims %left_idx_335 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc170)
+    %left_idx_337 = tt.broadcast %left_idx_336 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc171)
+    %right_idx_338 = arith.muli %y_idx_333, %flip_102 : tensor<64x2x4xi32, #linear3> loc(#loc173)
+    %right_idx_339 = "tt.reduce"(%right_idx_338) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<64x2x4xi32, #linear3>) -> tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> loc(#loc201)
+    %right_idx_340 = tt.expand_dims %right_idx_339 {axis = 1 : i32} : tensor<64x4xi32, #ttg.slice<{dim = 1, parent = #linear3}>> -> tensor<64x1x4xi32, #linear3> loc(#loc175)
+    %right_idx_341 = tt.broadcast %right_idx_340 : tensor<64x1x4xi32, #linear3> -> tensor<64x2x4xi32, #linear3> loc(#loc176)
+    %left_idx_342 = tt.reshape %left_idx_337 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_343 = tt.reshape %right_idx_341 : tensor<64x2x4xi32, #linear3> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_344 = arith.cmpi slt, %ileft_331, %iright_332 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_345 = arith.cmpi eq, %ileft_331, %iright_332 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_346 = arith.cmpi sgt, %left_idx_342, %right_idx_343 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_347 = arith.andi %eq_345, %cond_346 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_348 = arith.ori %cond_344, %cond_347 : tensor<32x16xi1, #linear> loc(#loc183)
+    %ret_349 = arith.xori %ileft_331, %iright_332 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_350 = arith.select %cond_348, %ret_349, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_351 = arith.xori %ret_318, %ret_350 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_352 = arith.xori %left_idx_342, %right_idx_343 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_353 = arith.select %cond_348, %new_idxs_352, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_354 = arith.xori %new_idxs_321, %new_idxs_353 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_355 = tt.reshape %ret_351 : tensor<32x16xi32, #linear> -> tensor<128x2x2xi32, #linear2> loc(#loc154)
+    %ileft_356 = arith.muli %y_355, %ileft_105 : tensor<128x2x2xi32, #linear2> loc(#loc156)
+    %ileft_357 = "tt.reduce"(%ileft_356) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc193)
+    %ileft_358 = tt.expand_dims %ileft_357 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc158)
+    %ileft_359 = tt.broadcast %ileft_358 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc159)
+    %iright_360 = arith.muli %y_355, %flip_48 : tensor<128x2x2xi32, #linear2> loc(#loc160)
+    %iright_361 = "tt.reduce"(%iright_360) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc195)
+    %iright_362 = tt.expand_dims %iright_361 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc162)
+    %iright_363 = tt.broadcast %iright_362 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc163)
+    %ileft_364 = tt.reshape %ileft_359 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_365 = tt.reshape %iright_363 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_366 = tt.reshape %new_idxs_354 : tensor<32x16xi32, #linear> -> tensor<128x2x2xi32, #linear2> loc(#loc166)
+    %left_idx_367 = arith.muli %y_idx_366, %ileft_105 : tensor<128x2x2xi32, #linear2> loc(#loc168)
+    %left_idx_368 = "tt.reduce"(%left_idx_367) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc198)
+    %left_idx_369 = tt.expand_dims %left_idx_368 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc170)
+    %left_idx_370 = tt.broadcast %left_idx_369 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc171)
+    %right_idx_371 = arith.muli %y_idx_366, %flip_48 : tensor<128x2x2xi32, #linear2> loc(#loc173)
+    %right_idx_372 = "tt.reduce"(%right_idx_371) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<128x2x2xi32, #linear2>) -> tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> loc(#loc201)
+    %right_idx_373 = tt.expand_dims %right_idx_372 {axis = 1 : i32} : tensor<128x2xi32, #ttg.slice<{dim = 1, parent = #linear2}>> -> tensor<128x1x2xi32, #linear2> loc(#loc175)
+    %right_idx_374 = tt.broadcast %right_idx_373 : tensor<128x1x2xi32, #linear2> -> tensor<128x2x2xi32, #linear2> loc(#loc176)
+    %left_idx_375 = tt.reshape %left_idx_370 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_376 = tt.reshape %right_idx_374 : tensor<128x2x2xi32, #linear2> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_377 = arith.cmpi slt, %ileft_364, %iright_365 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_378 = arith.cmpi eq, %ileft_364, %iright_365 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_379 = arith.cmpi sgt, %left_idx_375, %right_idx_376 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_380 = arith.andi %eq_378, %cond_379 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_381 = arith.ori %cond_377, %cond_380 : tensor<32x16xi1, #linear> loc(#loc183)
+    %ret_382 = arith.xori %ileft_364, %iright_365 : tensor<32x16xi32, #linear> loc(#loc186)
+    %ret_383 = arith.select %cond_381, %ret_382, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc187)
+    %ret_384 = arith.xori %ret_351, %ret_383 : tensor<32x16xi32, #linear> loc(#loc188)
+    %new_idxs_385 = arith.xori %left_idx_375, %right_idx_376 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_386 = arith.select %cond_381, %new_idxs_385, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_387 = arith.xori %new_idxs_354, %new_idxs_386 : tensor<32x16xi32, #linear> loc(#loc191)
+    %y_388 = tt.reshape %ret_384 : tensor<32x16xi32, #linear> -> tensor<256x2x1xi32, #linear1> loc(#loc154)
+    %ileft_389 = arith.muli %y_388, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc156)
+    %ileft_390 = "tt.reduce"(%ileft_389) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_419: i32 loc(callsite(#loc1 at #loc157)), %ileft_420: i32 loc(callsite(#loc1 at #loc157))):
+      %ileft_421 = arith.addi %ileft_419, %ileft_420 : i32 loc(#loc203)
+      tt.reduce.return %ileft_421 : i32 loc(#loc193)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc193)
+    %ileft_391 = tt.expand_dims %ileft_390 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc158)
+    %ileft_392 = tt.broadcast %ileft_391 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc159)
+    %iright_393 = arith.muli %y_388, %iright : tensor<256x2x1xi32, #linear1> loc(#loc160)
+    %iright_394 = "tt.reduce"(%iright_393) <{axis = 1 : i32}> ({
+    ^bb0(%iright_419: i32 loc(callsite(#loc1 at #loc161)), %iright_420: i32 loc(callsite(#loc1 at #loc161))):
+      %iright_421 = arith.addi %iright_419, %iright_420 : i32 loc(#loc204)
+      tt.reduce.return %iright_421 : i32 loc(#loc195)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc195)
+    %iright_395 = tt.expand_dims %iright_394 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc162)
+    %iright_396 = tt.broadcast %iright_395 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc163)
+    %ileft_397 = tt.reshape %ileft_392 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc164)
+    %iright_398 = tt.reshape %iright_396 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc165)
+    %y_idx_399 = tt.reshape %new_idxs_387 : tensor<32x16xi32, #linear> -> tensor<256x2x1xi32, #linear1> loc(#loc166)
+    %left_idx_400 = arith.muli %y_idx_399, %ileft : tensor<256x2x1xi32, #linear1> loc(#loc168)
+    %left_idx_401 = "tt.reduce"(%left_idx_400) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_419: i32 loc(callsite(#loc1 at #loc169)), %left_idx_420: i32 loc(callsite(#loc1 at #loc169))):
+      %left_idx_421 = arith.addi %left_idx_419, %left_idx_420 : i32 loc(#loc205)
+      tt.reduce.return %left_idx_421 : i32 loc(#loc198)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc198)
+    %left_idx_402 = tt.expand_dims %left_idx_401 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc170)
+    %left_idx_403 = tt.broadcast %left_idx_402 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc171)
+    %right_idx_404 = arith.muli %y_idx_399, %iright : tensor<256x2x1xi32, #linear1> loc(#loc173)
+    %right_idx_405 = "tt.reduce"(%right_idx_404) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_419: i32 loc(callsite(#loc1 at #loc174)), %right_idx_420: i32 loc(callsite(#loc1 at #loc174))):
+      %right_idx_421 = arith.addi %right_idx_419, %right_idx_420 : i32 loc(#loc206)
+      tt.reduce.return %right_idx_421 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32, #linear1>) -> tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> loc(#loc201)
+    %right_idx_406 = tt.expand_dims %right_idx_405 {axis = 1 : i32} : tensor<256x1xi32, #ttg.slice<{dim = 1, parent = #linear1}>> -> tensor<256x1x1xi32, #linear1> loc(#loc175)
+    %right_idx_407 = tt.broadcast %right_idx_406 : tensor<256x1x1xi32, #linear1> -> tensor<256x2x1xi32, #linear1> loc(#loc176)
+    %left_idx_408 = tt.reshape %left_idx_403 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc177)
+    %right_idx_409 = tt.reshape %right_idx_407 : tensor<256x2x1xi32, #linear1> -> tensor<32x16xi32, #linear> loc(#loc178)
+    %cond_410 = arith.cmpi slt, %ileft_397, %iright_398 : tensor<32x16xi32, #linear> loc(#loc179)
+    %eq_411 = arith.cmpi eq, %ileft_397, %iright_398 : tensor<32x16xi32, #linear> loc(#loc180)
+    %cond_412 = arith.cmpi sgt, %left_idx_408, %right_idx_409 : tensor<32x16xi32, #linear> loc(#loc181)
+    %cond_413 = arith.andi %eq_411, %cond_412 : tensor<32x16xi1, #linear> loc(#loc182)
+    %cond_414 = arith.ori %cond_410, %cond_413 : tensor<32x16xi1, #linear> loc(#loc183)
+    %new_idxs_415 = arith.xori %left_idx_408, %right_idx_409 : tensor<32x16xi32, #linear> loc(#loc189)
+    %new_idxs_416 = arith.select %cond_414, %new_idxs_415, %cst : tensor<32x16xi1, #linear>, tensor<32x16xi32, #linear> loc(#loc190)
+    %new_idxs_417 = arith.xori %new_idxs_387, %new_idxs_416 : tensor<32x16xi32, #linear> loc(#loc191)
+    %tmp7 = arith.extsi %tmp0_36 : tensor<32x16xi32, #blocked> to tensor<32x16xi64, #blocked> loc(#loc141)
+    %tmp10 = arith.select %tmp0_34, %tmp7, %cst_0 : tensor<32x16xi1, #blocked>, tensor<32x16xi64, #blocked> loc(#loc142)
+    %tmp11 = "tt.reduce"(%tmp10) <{axis = 1 : i32}> ({
+    ^bb0(%tmp11_419: i64 loc(callsite(#loc1 at #loc143)), %tmp11_420: i64 loc(callsite(#loc1 at #loc143))):
+      %tmp11_421 = arith.addi %tmp11_419, %tmp11_420 : i64 loc(#loc192)
+      tt.reduce.return %tmp11_421 : i64 loc(#loc152)
+    }) : (tensor<32x16xi64, #blocked>) -> tensor<32xi64, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc152)
+    %tmp11_418 = tt.expand_dims %tmp11 {axis = 1 : i32} : tensor<32xi64, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi64, #blocked> loc(#loc144)
+    %tmp14 = arith.trunci %tmp11_418 : tensor<32x1xi64, #blocked> to tensor<32x1xi32, #blocked> loc(#loc145)
+    %0 = arith.muli %xindex_19, %cst_4 : tensor<32x1xi32, #blocked1> loc(#loc70)
+    %1 = tt.broadcast %r0_index_25 : tensor<1x16xi32, #blocked1> -> tensor<32x16xi32, #blocked1> loc(#loc71)
+    %2 = tt.broadcast %0 : tensor<32x1xi32, #blocked1> -> tensor<32x16xi32, #blocked1> loc(#loc71)
+    %3 = arith.addi %1, %2 : tensor<32x16xi32, #blocked1> loc(#loc71)
+    %4 = tt.splat %out_ptr2 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>, #blocked1> loc(#loc72)
+    %5 = tt.addptr %4, %3 : tensor<32x16x!tt.ptr<i32>, #blocked1>, tensor<32x16xi32, #blocked1> loc(#loc72)
+    %6 = ttg.convert_layout %new_idxs_417 : tensor<32x16xi32, #linear> -> tensor<32x16xi32, #blocked1> loc(#loc73)
+    tt.store %5, %6, %tmp0_35 : tensor<32x16x!tt.ptr<i32>, #blocked1> loc(#loc73)
+    %7 = tt.splat %out_ptr3 : !tt.ptr<i32> -> tensor<32x1x!tt.ptr<i32>, #blocked> loc(#loc74)
+    %8 = tt.addptr %7, %xindex_18 : tensor<32x1x!tt.ptr<i32>, #blocked>, tensor<32x1xi32, #blocked> loc(#loc74)
+    tt.store %8, %tmp14, %xmask : tensor<32x1x!tt.ptr<i32>, #blocked> loc(#loc75)
+    tt.return loc(#loc76)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":24:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":24:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":25:44)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":25:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":26:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":27:38)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":33:19)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":34:19)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:38)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:35)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:49)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:45)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:30)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:54)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":38:19)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":40:33)
+#loc18 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:44)
+#loc21 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:60)
+#loc22 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:68)
+#loc23 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":533:22)
+#loc25 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":537:21)
+#loc26 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:40)
+#loc27 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc29 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:65)
+#loc31 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:78)
+#loc32 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:41)
+#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:67)
+#loc35 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:80)
+#loc36 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":540:30)
+#loc37 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":541:32)
+#loc38 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":546:29)
+#loc39 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:36)
+#loc40 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:23)
+#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":290:25)
+#loc43 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:53)
+#loc44 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:66)
+#loc45 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:37)
+#loc46 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:23)
+#loc48 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:54)
+#loc49 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:67)
+#loc50 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":553:36)
+#loc51 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":554:38)
+#loc52 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":574:22)
+#loc53 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":591:21)
+#loc54 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:40)
+#loc55 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:29)
+#loc56 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:23)
+#loc57 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":599:19)
+#loc58 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":599:28)
+#loc59 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:38)
+#loc60 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:46)
+#loc61 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:15)
+#loc62 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:48)
+#loc63 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:59)
+#loc64 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:22)
+#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":42:19)
+#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":44:34)
+#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":45:29)
+#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":48:21)
+#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:35)
+#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:32)
+#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:25)
+#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:47)
+#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":50:25)
+#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":50:37)
+#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":50:4)
+#loc82 = loc("xoffset"(#loc2))
+#loc83 = loc("xoffset"(#loc3))
+#loc84 = loc("xindex"(#loc4))
+#loc85 = loc("xindex"(#loc5))
+#loc86 = loc("xmask"(#loc6))
+#loc87 = loc("r0_index"(#loc7))
+#loc88 = loc("x0"(#loc8))
+#loc89 = loc("x1"(#loc9))
+#loc90 = loc("tmp0"(#loc10))
+#loc91 = loc("tmp0"(#loc11))
+#loc92 = loc("tmp0"(#loc12))
+#loc93 = loc("tmp0"(#loc13))
+#loc94 = loc("tmp0"(#loc14))
+#loc95 = loc("tmp0"(#loc15))
+#loc96 = loc("tmp2"(#loc16))
+#loc97 = loc("tmp4"(#loc17))
+#loc98 = loc("flip"(#loc18))
+#loc100 = loc("flip"(#loc21))
+#loc101 = loc("flip"(#loc22))
+#loc102 = loc("y"(#loc23))
+#loc103 = loc("left_mask"(#loc25))
+#loc104 = loc("ileft"(#loc26))
+#loc106 = loc("ileft"(#loc30))
+#loc107 = loc("ileft"(#loc31))
+#loc108 = loc("iright"(#loc32))
+#loc110 = loc("iright"(#loc34))
+#loc111 = loc("iright"(#loc35))
+#loc112 = loc("ileft"(#loc36))
+#loc113 = loc("iright"(#loc37))
+#loc114 = loc("y_idx"(#loc38))
+#loc115 = loc("left_idx"(#loc39))
+#loc116 = loc("left_idx"(#loc40))
+#loc117 = loc("input"(#loc41))
+#loc119 = loc("left_idx"(#loc43))
+#loc120 = loc("left_idx"(#loc44))
+#loc121 = loc("right_idx"(#loc45))
+#loc122 = loc("right_idx"(#loc46))
+#loc124 = loc("right_idx"(#loc48))
+#loc125 = loc("right_idx"(#loc49))
+#loc126 = loc("left_idx"(#loc50))
+#loc127 = loc("right_idx"(#loc51))
+#loc128 = loc("cond"(#loc52))
+#loc129 = loc("eq"(#loc53))
+#loc130 = loc("cond"(#loc54))
+#loc131 = loc("cond"(#loc55))
+#loc132 = loc("cond"(#loc56))
+#loc133 = loc("cond"(#loc57))
+#loc134 = loc("cond"(#loc58))
+#loc135 = loc("ret"(#loc59))
+#loc136 = loc("ret"(#loc60))
+#loc137 = loc("ret"(#loc61))
+#loc138 = loc("new_idxs"(#loc62))
+#loc139 = loc("new_idxs"(#loc63))
+#loc140 = loc("new_idxs"(#loc64))
+#loc141 = loc("tmp7"(#loc65))
+#loc142 = loc("tmp10"(#loc66))
+#loc144 = loc("tmp11"(#loc68))
+#loc145 = loc("tmp14"(#loc69))
+#loc146 = loc(callsite(#loc98 at #loc99))
+#loc147 = loc(callsite(#loc100 at #loc99))
+#loc148 = loc(callsite(#loc101 at #loc99))
+#loc150 = loc("cond"(#loc128))
+#loc151 = loc("eq"(#loc129))
+#loc152 = loc(callsite(#loc27 at #loc143))
+#loc154 = loc(callsite(#loc102 at #loc149))
+#loc155 = loc(callsite(#loc103 at #loc149))
+#loc156 = loc(callsite(#loc104 at #loc149))
+#loc158 = loc(callsite(#loc106 at #loc149))
+#loc159 = loc(callsite(#loc107 at #loc149))
+#loc160 = loc(callsite(#loc108 at #loc149))
+#loc162 = loc(callsite(#loc110 at #loc149))
+#loc163 = loc(callsite(#loc111 at #loc149))
+#loc164 = loc(callsite(#loc112 at #loc149))
+#loc165 = loc(callsite(#loc113 at #loc149))
+#loc166 = loc(callsite(#loc114 at #loc149))
+#loc167 = loc(callsite(#loc115 at #loc149))
+#loc168 = loc(callsite(#loc116 at #loc149))
+#loc170 = loc(callsite(#loc119 at #loc149))
+#loc171 = loc(callsite(#loc120 at #loc149))
+#loc172 = loc(callsite(#loc121 at #loc149))
+#loc173 = loc(callsite(#loc122 at #loc149))
+#loc175 = loc(callsite(#loc124 at #loc149))
+#loc176 = loc(callsite(#loc125 at #loc149))
+#loc177 = loc(callsite(#loc126 at #loc149))
+#loc178 = loc(callsite(#loc127 at #loc149))
+#loc179 = loc(callsite(#loc150 at #loc149))
+#loc180 = loc(callsite(#loc151 at #loc149))
+#loc181 = loc(callsite(#loc130 at #loc149))
+#loc182 = loc(callsite(#loc131 at #loc149))
+#loc183 = loc(callsite(#loc132 at #loc149))
+#loc184 = loc(callsite(#loc133 at #loc149))
+#loc185 = loc(callsite(#loc134 at #loc149))
+#loc186 = loc(callsite(#loc135 at #loc149))
+#loc187 = loc(callsite(#loc136 at #loc149))
+#loc188 = loc(callsite(#loc137 at #loc149))
+#loc189 = loc(callsite(#loc138 at #loc149))
+#loc190 = loc(callsite(#loc139 at #loc149))
+#loc191 = loc(callsite(#loc140 at #loc149))
+#loc192 = loc(callsite(#loc29 at #loc152))
+#loc193 = loc(callsite(#loc27 at #loc157))
+#loc195 = loc(callsite(#loc27 at #loc161))
+#loc197 = loc(callsite(#loc117 at #loc169))
+#loc198 = loc(callsite(#loc27 at #loc169))
+#loc200 = loc(callsite(#loc117 at #loc174))
+#loc201 = loc(callsite(#loc27 at #loc174))
+#loc203 = loc(callsite(#loc29 at #loc193))
+#loc204 = loc(callsite(#loc29 at #loc195))
+#loc205 = loc(callsite(#loc29 at #loc198))
+#loc206 = loc(callsite(#loc29 at #loc201))

SpecForge-ext/cache/compiled_kernels/triton/6/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttir

@@ -0,0 +1,799 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":18:0)
+#loc1 = loc(unknown)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":41:67)
+#loc23 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":662:12)
+#loc28 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":634:73)
+#loc32 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:51)
+#loc37 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:53)
+#loc46 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:50)
+#loc51 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:51)
+#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":45:26)
+#loc80 = loc("in_ptr0"(#loc))
+#loc81 = loc("out_ptr2"(#loc))
+#loc82 = loc("out_ptr3"(#loc))
+#loc83 = loc("xnumel"(#loc))
+#loc84 = loc("r0_numel"(#loc))
+#loc106 = loc(callsite(#loc23 at #loc2))
+#loc113 = loc("ileft"(#loc32))
+#loc117 = loc("iright"(#loc37))
+#loc126 = loc("left_idx"(#loc46))
+#loc131 = loc("right_idx"(#loc51))
+#loc150 = loc("tmp11"(#loc70))
+#loc157 = loc(callsite(#loc28 at #loc106))
+#loc161 = loc(callsite(#loc1 at #loc150))
+#loc165 = loc(callsite(#loc113 at #loc157))
+#loc169 = loc(callsite(#loc117 at #loc157))
+#loc177 = loc(callsite(#loc126 at #loc157))
+#loc182 = loc(callsite(#loc131 at #loc157))
+#loc202 = loc(callsite(#loc1 at #loc165))
+#loc204 = loc(callsite(#loc1 at #loc169))
+#loc207 = loc(callsite(#loc1 at #loc177))
+#loc210 = loc(callsite(#loc1 at #loc182))
+module {
+  tt.func public @triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3(%in_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr2: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr2"(#loc)), %out_ptr3: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr3"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<1> : tensor<1x2x1xi32> loc(#loc85)
+    %cst_0 = arith.constant dense<0> : tensor<32x16xi32> loc(#loc1)
+    %tmp10 = arith.constant dense<0> : tensor<32x16xi64> loc(#loc86)
+    %tmp0 = arith.constant dense<272> : tensor<32x1xi32> loc(#loc87)
+    %tmp0_1 = arith.constant dense<17> : tensor<1x16xi32> loc(#loc88)
+    %cst_2 = arith.constant dense<16> : tensor<32x1xi32> loc(#loc1)
+    %xmask = arith.constant dense<128> : tensor<32x1xi32> loc(#loc89)
+    %c32_i32 = arith.constant 32 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc90)
+    %xoffset_3 = arith.muli %xoffset, %c32_i32 : i32 loc(#loc91)
+    %xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc92)
+    %xindex_4 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc93)
+    %xindex_5 = tt.splat %xoffset_3 : i32 -> tensor<32x1xi32> loc(#loc94)
+    %xindex_6 = arith.addi %xindex_5, %xindex_4 : tensor<32x1xi32> loc(#loc94)
+    %xmask_7 = arith.cmpi slt, %xindex_6, %xmask : tensor<32x1xi32> loc(#loc89)
+    %r0_index = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc95)
+    %r0_index_8 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc96)
+    %x0 = arith.remsi %xindex_6, %cst_2 : tensor<32x1xi32> loc(#loc97)
+    %x1 = arith.divsi %xindex_6, %cst_2 : tensor<32x1xi32> loc(#loc98)
+    %tmp0_9 = arith.muli %r0_index_8, %tmp0_1 : tensor<1x16xi32> loc(#loc88)
+    %tmp0_10 = tt.broadcast %x0 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc99)
+    %tmp0_11 = tt.broadcast %tmp0_9 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc99)
+    %tmp0_12 = arith.addi %tmp0_10, %tmp0_11 : tensor<32x16xi32> loc(#loc99)
+    %tmp0_13 = arith.muli %x1, %tmp0 : tensor<32x1xi32> loc(#loc87)
+    %tmp0_14 = tt.broadcast %tmp0_13 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc100)
+    %tmp0_15 = arith.addi %tmp0_12, %tmp0_14 : tensor<32x16xi32> loc(#loc100)
+    %tmp0_16 = tt.splat %in_ptr0 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>> loc(#loc101)
+    %tmp0_17 = tt.addptr %tmp0_16, %tmp0_15 : tensor<32x16x!tt.ptr<i32>>, tensor<32x16xi32> loc(#loc101)
+    %tmp0_18 = tt.broadcast %xmask_7 : tensor<32x1xi1> -> tensor<32x16xi1> loc(#loc102)
+    %tmp0_19 = tt.load %tmp0_17, %tmp0_18, %cst_0 : tensor<32x16x!tt.ptr<i32>> loc(#loc102)
+    %tmp2 = arith.trunci %r0_index_8 : tensor<1x16xi32> to tensor<1x16xi16> loc(#loc103)
+    %tmp4 = tt.broadcast %tmp2 : tensor<1x16xi16> -> tensor<32x16xi16> loc(#loc104)
+    %flip = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32> loc(#loc153)
+    %flip_20 = tt.expand_dims %flip {axis = 0 : i32} : tensor<2xi32> -> tensor<1x2xi32> loc(#loc154)
+    %flip_21 = tt.expand_dims %flip_20 {axis = 2 : i32} : tensor<1x2xi32> -> tensor<1x2x1xi32> loc(#loc154)
+    %flip_22 = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<128x2x2xi32> loc(#loc155)
+    %flip_23 = tt.reshape %flip_22 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc156)
+    %y = tt.reshape %tmp0_19 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
+    %left_mask = arith.subi %cst, %flip_21 : tensor<1x2x1xi32> loc(#loc163)
+    %ileft = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<256x2x1xi32> loc(#loc164)
+    %ileft_24 = arith.muli %y, %ileft : tensor<256x2x1xi32> loc(#loc164)
+    %ileft_25 = "tt.reduce"(%ileft_24) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
+    %ileft_26 = tt.expand_dims %ileft_25 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
+    %ileft_27 = tt.broadcast %ileft_26 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
+    %iright = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<256x2x1xi32> loc(#loc168)
+    %iright_28 = arith.muli %y, %iright : tensor<256x2x1xi32> loc(#loc168)
+    %iright_29 = "tt.reduce"(%iright_28) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
+    %iright_30 = tt.expand_dims %iright_29 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
+    %iright_31 = tt.broadcast %iright_30 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
+    %ileft_32 = tt.reshape %ileft_27 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_33 = tt.reshape %iright_31 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx = tt.reshape %tmp4 : tensor<32x16xi16> -> tensor<256x2x1xi16> loc(#loc174)
+    %left_idx = arith.trunci %left_mask : tensor<1x2x1xi32> to tensor<1x2x1xi16> loc(#loc175)
+    %left_idx_34 = tt.broadcast %left_idx : tensor<1x2x1xi16> -> tensor<256x2x1xi16> loc(#loc176)
+    %left_idx_35 = arith.muli %y_idx, %left_idx_34 : tensor<256x2x1xi16> loc(#loc176)
+    %input = arith.extsi %left_idx_35 : tensor<256x2x1xi16> to tensor<256x2x1xi32> loc(#loc205)
+    %left_idx_36 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
+    %left_idx_37 = tt.expand_dims %left_idx_36 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
+    %left_idx_38 = tt.broadcast %left_idx_37 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
+    %right_idx = arith.trunci %flip_21 : tensor<1x2x1xi32> to tensor<1x2x1xi16> loc(#loc180)
+    %right_idx_39 = tt.broadcast %right_idx : tensor<1x2x1xi16> -> tensor<256x2x1xi16> loc(#loc181)
+    %right_idx_40 = arith.muli %y_idx, %right_idx_39 : tensor<256x2x1xi16> loc(#loc181)
+    %input_41 = arith.extsi %right_idx_40 : tensor<256x2x1xi16> to tensor<256x2x1xi32> loc(#loc208)
+    %right_idx_42 = "tt.reduce"(%input_41) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
+    %right_idx_43 = tt.expand_dims %right_idx_42 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
+    %right_idx_44 = tt.broadcast %right_idx_43 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
+    %left_idx_45 = tt.reshape %left_idx_38 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_46 = tt.reshape %right_idx_44 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond = arith.cmpi slt, %ileft_32, %iright_33 : tensor<32x16xi32> loc(#loc187)
+    %eq = arith.cmpi eq, %ileft_32, %iright_33 : tensor<32x16xi32> loc(#loc188)
+    %cond_47 = arith.cmpi sgt, %left_idx_45, %right_idx_46 : tensor<32x16xi32> loc(#loc189)
+    %cond_48 = arith.andi %eq, %cond_47 : tensor<32x16xi1> loc(#loc190)
+    %cond_49 = arith.ori %cond, %cond_48 : tensor<32x16xi1> loc(#loc191)
+    %cond_50 = arith.extui %cond_49 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
+    %cond_51 = arith.xori %cond_50, %flip_23 : tensor<32x16xi32> loc(#loc192)
+    %cond_52 = arith.cmpi ne, %cond_51, %cst_0 : tensor<32x16xi32> loc(#loc193)
+    %ret = arith.xori %ileft_32, %iright_33 : tensor<32x16xi32> loc(#loc194)
+    %ret_53 = arith.select %cond_52, %ret, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_54 = arith.xori %tmp0_19, %ret_53 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs = arith.xori %left_idx_45, %right_idx_46 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_55 = arith.select %cond_52, %new_idxs, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_56 = arith.extsi %tmp2 : tensor<1x16xi16> to tensor<1x16xi32> loc(#loc199)
+    %new_idxs_57 = tt.broadcast %new_idxs_56 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc199)
+    %new_idxs_58 = arith.xori %new_idxs_57, %new_idxs_55 : tensor<32x16xi32> loc(#loc199)
+    %flip_59 = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<64x2x4xi32> loc(#loc155)
+    %flip_60 = tt.reshape %flip_59 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc156)
+    %y_61 = tt.reshape %ret_54 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc162)
+    %ileft_62 = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<128x2x2xi32> loc(#loc164)
+    %ileft_63 = arith.muli %y_61, %ileft_62 : tensor<128x2x2xi32> loc(#loc164)
+    %ileft_64 = "tt.reduce"(%ileft_63) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc201)
+    %ileft_65 = tt.expand_dims %ileft_64 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc166)
+    %ileft_66 = tt.broadcast %ileft_65 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc167)
+    %iright_67 = arith.muli %y_61, %flip_22 : tensor<128x2x2xi32> loc(#loc168)
+    %iright_68 = "tt.reduce"(%iright_67) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc203)
+    %iright_69 = tt.expand_dims %iright_68 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc170)
+    %iright_70 = tt.broadcast %iright_69 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc171)
+    %ileft_71 = tt.reshape %ileft_66 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_72 = tt.reshape %iright_70 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_73 = tt.reshape %new_idxs_58 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc174)
+    %left_idx_74 = arith.muli %y_idx_73, %ileft_62 : tensor<128x2x2xi32> loc(#loc176)
+    %left_idx_75 = "tt.reduce"(%left_idx_74) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc206)
+    %left_idx_76 = tt.expand_dims %left_idx_75 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc178)
+    %left_idx_77 = tt.broadcast %left_idx_76 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc179)
+    %right_idx_78 = arith.muli %y_idx_73, %flip_22 : tensor<128x2x2xi32> loc(#loc181)
+    %right_idx_79 = "tt.reduce"(%right_idx_78) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc209)
+    %right_idx_80 = tt.expand_dims %right_idx_79 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc183)
+    %right_idx_81 = tt.broadcast %right_idx_80 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc184)
+    %left_idx_82 = tt.reshape %left_idx_77 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_83 = tt.reshape %right_idx_81 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_84 = arith.cmpi slt, %ileft_71, %iright_72 : tensor<32x16xi32> loc(#loc187)
+    %eq_85 = arith.cmpi eq, %ileft_71, %iright_72 : tensor<32x16xi32> loc(#loc188)
+    %cond_86 = arith.cmpi sgt, %left_idx_82, %right_idx_83 : tensor<32x16xi32> loc(#loc189)
+    %cond_87 = arith.andi %eq_85, %cond_86 : tensor<32x16xi1> loc(#loc190)
+    %cond_88 = arith.ori %cond_84, %cond_87 : tensor<32x16xi1> loc(#loc191)
+    %cond_89 = arith.extui %cond_88 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
+    %cond_90 = arith.xori %cond_89, %flip_60 : tensor<32x16xi32> loc(#loc192)
+    %cond_91 = arith.cmpi ne, %cond_90, %cst_0 : tensor<32x16xi32> loc(#loc193)
+    %ret_92 = arith.xori %ileft_71, %iright_72 : tensor<32x16xi32> loc(#loc194)
+    %ret_93 = arith.select %cond_91, %ret_92, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_94 = arith.xori %ret_54, %ret_93 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_95 = arith.xori %left_idx_82, %right_idx_83 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_96 = arith.select %cond_91, %new_idxs_95, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_97 = arith.xori %new_idxs_58, %new_idxs_96 : tensor<32x16xi32> loc(#loc199)
+    %y_98 = tt.reshape %ret_94 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
+    %ileft_99 = arith.muli %y_98, %ileft : tensor<256x2x1xi32> loc(#loc164)
+    %ileft_100 = "tt.reduce"(%ileft_99) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
+    %ileft_101 = tt.expand_dims %ileft_100 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
+    %ileft_102 = tt.broadcast %ileft_101 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
+    %iright_103 = arith.muli %y_98, %iright : tensor<256x2x1xi32> loc(#loc168)
+    %iright_104 = "tt.reduce"(%iright_103) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
+    %iright_105 = tt.expand_dims %iright_104 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
+    %iright_106 = tt.broadcast %iright_105 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
+    %ileft_107 = tt.reshape %ileft_102 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_108 = tt.reshape %iright_106 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_109 = tt.reshape %new_idxs_97 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc174)
+    %left_idx_110 = arith.muli %y_idx_109, %ileft : tensor<256x2x1xi32> loc(#loc176)
+    %left_idx_111 = "tt.reduce"(%left_idx_110) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
+    %left_idx_112 = tt.expand_dims %left_idx_111 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
+    %left_idx_113 = tt.broadcast %left_idx_112 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
+    %right_idx_114 = arith.muli %y_idx_109, %iright : tensor<256x2x1xi32> loc(#loc181)
+    %right_idx_115 = "tt.reduce"(%right_idx_114) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
+    %right_idx_116 = tt.expand_dims %right_idx_115 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
+    %right_idx_117 = tt.broadcast %right_idx_116 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
+    %left_idx_118 = tt.reshape %left_idx_113 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_119 = tt.reshape %right_idx_117 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_120 = arith.cmpi slt, %ileft_107, %iright_108 : tensor<32x16xi32> loc(#loc187)
+    %eq_121 = arith.cmpi eq, %ileft_107, %iright_108 : tensor<32x16xi32> loc(#loc188)
+    %cond_122 = arith.cmpi sgt, %left_idx_118, %right_idx_119 : tensor<32x16xi32> loc(#loc189)
+    %cond_123 = arith.andi %eq_121, %cond_122 : tensor<32x16xi1> loc(#loc190)
+    %cond_124 = arith.ori %cond_120, %cond_123 : tensor<32x16xi1> loc(#loc191)
+    %cond_125 = arith.extui %cond_124 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
+    %cond_126 = arith.xori %cond_125, %flip_60 : tensor<32x16xi32> loc(#loc192)
+    %cond_127 = arith.cmpi ne, %cond_126, %cst_0 : tensor<32x16xi32> loc(#loc193)
+    %ret_128 = arith.xori %ileft_107, %iright_108 : tensor<32x16xi32> loc(#loc194)
+    %ret_129 = arith.select %cond_127, %ret_128, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_130 = arith.xori %ret_94, %ret_129 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_131 = arith.xori %left_idx_118, %right_idx_119 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_132 = arith.select %cond_127, %new_idxs_131, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_133 = arith.xori %new_idxs_97, %new_idxs_132 : tensor<32x16xi32> loc(#loc199)
+    %flip_134 = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<32x2x8xi32> loc(#loc155)
+    %flip_135 = tt.reshape %flip_134 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc156)
+    %y_136 = tt.reshape %ret_130 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc162)
+    %ileft_137 = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<64x2x4xi32> loc(#loc164)
+    %ileft_138 = arith.muli %y_136, %ileft_137 : tensor<64x2x4xi32> loc(#loc164)
+    %ileft_139 = "tt.reduce"(%ileft_138) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc201)
+    %ileft_140 = tt.expand_dims %ileft_139 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc166)
+    %ileft_141 = tt.broadcast %ileft_140 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc167)
+    %iright_142 = arith.muli %y_136, %flip_59 : tensor<64x2x4xi32> loc(#loc168)
+    %iright_143 = "tt.reduce"(%iright_142) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc203)
+    %iright_144 = tt.expand_dims %iright_143 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc170)
+    %iright_145 = tt.broadcast %iright_144 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc171)
+    %ileft_146 = tt.reshape %ileft_141 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_147 = tt.reshape %iright_145 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_148 = tt.reshape %new_idxs_133 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc174)
+    %left_idx_149 = arith.muli %y_idx_148, %ileft_137 : tensor<64x2x4xi32> loc(#loc176)
+    %left_idx_150 = "tt.reduce"(%left_idx_149) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc206)
+    %left_idx_151 = tt.expand_dims %left_idx_150 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc178)
+    %left_idx_152 = tt.broadcast %left_idx_151 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc179)
+    %right_idx_153 = arith.muli %y_idx_148, %flip_59 : tensor<64x2x4xi32> loc(#loc181)
+    %right_idx_154 = "tt.reduce"(%right_idx_153) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc209)
+    %right_idx_155 = tt.expand_dims %right_idx_154 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc183)
+    %right_idx_156 = tt.broadcast %right_idx_155 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc184)
+    %left_idx_157 = tt.reshape %left_idx_152 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_158 = tt.reshape %right_idx_156 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_159 = arith.cmpi slt, %ileft_146, %iright_147 : tensor<32x16xi32> loc(#loc187)
+    %eq_160 = arith.cmpi eq, %ileft_146, %iright_147 : tensor<32x16xi32> loc(#loc188)
+    %cond_161 = arith.cmpi sgt, %left_idx_157, %right_idx_158 : tensor<32x16xi32> loc(#loc189)
+    %cond_162 = arith.andi %eq_160, %cond_161 : tensor<32x16xi1> loc(#loc190)
+    %cond_163 = arith.ori %cond_159, %cond_162 : tensor<32x16xi1> loc(#loc191)
+    %cond_164 = arith.extui %cond_163 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
+    %cond_165 = arith.xori %cond_164, %flip_135 : tensor<32x16xi32> loc(#loc192)
+    %cond_166 = arith.cmpi ne, %cond_165, %cst_0 : tensor<32x16xi32> loc(#loc193)
+    %ret_167 = arith.xori %ileft_146, %iright_147 : tensor<32x16xi32> loc(#loc194)
+    %ret_168 = arith.select %cond_166, %ret_167, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_169 = arith.xori %ret_130, %ret_168 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_170 = arith.xori %left_idx_157, %right_idx_158 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_171 = arith.select %cond_166, %new_idxs_170, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_172 = arith.xori %new_idxs_133, %new_idxs_171 : tensor<32x16xi32> loc(#loc199)
+    %y_173 = tt.reshape %ret_169 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc162)
+    %ileft_174 = arith.muli %y_173, %ileft_62 : tensor<128x2x2xi32> loc(#loc164)
+    %ileft_175 = "tt.reduce"(%ileft_174) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc201)
+    %ileft_176 = tt.expand_dims %ileft_175 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc166)
+    %ileft_177 = tt.broadcast %ileft_176 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc167)
+    %iright_178 = arith.muli %y_173, %flip_22 : tensor<128x2x2xi32> loc(#loc168)
+    %iright_179 = "tt.reduce"(%iright_178) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc203)
+    %iright_180 = tt.expand_dims %iright_179 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc170)
+    %iright_181 = tt.broadcast %iright_180 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc171)
+    %ileft_182 = tt.reshape %ileft_177 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_183 = tt.reshape %iright_181 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_184 = tt.reshape %new_idxs_172 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc174)
+    %left_idx_185 = arith.muli %y_idx_184, %ileft_62 : tensor<128x2x2xi32> loc(#loc176)
+    %left_idx_186 = "tt.reduce"(%left_idx_185) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc206)
+    %left_idx_187 = tt.expand_dims %left_idx_186 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc178)
+    %left_idx_188 = tt.broadcast %left_idx_187 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc179)
+    %right_idx_189 = arith.muli %y_idx_184, %flip_22 : tensor<128x2x2xi32> loc(#loc181)
+    %right_idx_190 = "tt.reduce"(%right_idx_189) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc209)
+    %right_idx_191 = tt.expand_dims %right_idx_190 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc183)
+    %right_idx_192 = tt.broadcast %right_idx_191 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc184)
+    %left_idx_193 = tt.reshape %left_idx_188 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_194 = tt.reshape %right_idx_192 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_195 = arith.cmpi slt, %ileft_182, %iright_183 : tensor<32x16xi32> loc(#loc187)
+    %eq_196 = arith.cmpi eq, %ileft_182, %iright_183 : tensor<32x16xi32> loc(#loc188)
+    %cond_197 = arith.cmpi sgt, %left_idx_193, %right_idx_194 : tensor<32x16xi32> loc(#loc189)
+    %cond_198 = arith.andi %eq_196, %cond_197 : tensor<32x16xi1> loc(#loc190)
+    %cond_199 = arith.ori %cond_195, %cond_198 : tensor<32x16xi1> loc(#loc191)
+    %cond_200 = arith.extui %cond_199 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
+    %cond_201 = arith.xori %cond_200, %flip_135 : tensor<32x16xi32> loc(#loc192)
+    %cond_202 = arith.cmpi ne, %cond_201, %cst_0 : tensor<32x16xi32> loc(#loc193)
+    %ret_203 = arith.xori %ileft_182, %iright_183 : tensor<32x16xi32> loc(#loc194)
+    %ret_204 = arith.select %cond_202, %ret_203, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_205 = arith.xori %ret_169, %ret_204 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_206 = arith.xori %left_idx_193, %right_idx_194 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_207 = arith.select %cond_202, %new_idxs_206, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_208 = arith.xori %new_idxs_172, %new_idxs_207 : tensor<32x16xi32> loc(#loc199)
+    %y_209 = tt.reshape %ret_205 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
+    %ileft_210 = arith.muli %y_209, %ileft : tensor<256x2x1xi32> loc(#loc164)
+    %ileft_211 = "tt.reduce"(%ileft_210) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
+    %ileft_212 = tt.expand_dims %ileft_211 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
+    %ileft_213 = tt.broadcast %ileft_212 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
+    %iright_214 = arith.muli %y_209, %iright : tensor<256x2x1xi32> loc(#loc168)
+    %iright_215 = "tt.reduce"(%iright_214) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
+    %iright_216 = tt.expand_dims %iright_215 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
+    %iright_217 = tt.broadcast %iright_216 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
+    %ileft_218 = tt.reshape %ileft_213 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_219 = tt.reshape %iright_217 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_220 = tt.reshape %new_idxs_208 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc174)
+    %left_idx_221 = arith.muli %y_idx_220, %ileft : tensor<256x2x1xi32> loc(#loc176)
+    %left_idx_222 = "tt.reduce"(%left_idx_221) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
+    %left_idx_223 = tt.expand_dims %left_idx_222 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
+    %left_idx_224 = tt.broadcast %left_idx_223 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
+    %right_idx_225 = arith.muli %y_idx_220, %iright : tensor<256x2x1xi32> loc(#loc181)
+    %right_idx_226 = "tt.reduce"(%right_idx_225) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
+    %right_idx_227 = tt.expand_dims %right_idx_226 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
+    %right_idx_228 = tt.broadcast %right_idx_227 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
+    %left_idx_229 = tt.reshape %left_idx_224 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_230 = tt.reshape %right_idx_228 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_231 = arith.cmpi slt, %ileft_218, %iright_219 : tensor<32x16xi32> loc(#loc187)
+    %eq_232 = arith.cmpi eq, %ileft_218, %iright_219 : tensor<32x16xi32> loc(#loc188)
+    %cond_233 = arith.cmpi sgt, %left_idx_229, %right_idx_230 : tensor<32x16xi32> loc(#loc189)
+    %cond_234 = arith.andi %eq_232, %cond_233 : tensor<32x16xi1> loc(#loc190)
+    %cond_235 = arith.ori %cond_231, %cond_234 : tensor<32x16xi1> loc(#loc191)
+    %cond_236 = arith.extui %cond_235 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
+    %cond_237 = arith.xori %cond_236, %flip_135 : tensor<32x16xi32> loc(#loc192)
+    %cond_238 = arith.cmpi ne, %cond_237, %cst_0 : tensor<32x16xi32> loc(#loc193)
+    %ret_239 = arith.xori %ileft_218, %iright_219 : tensor<32x16xi32> loc(#loc194)
+    %ret_240 = arith.select %cond_238, %ret_239, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_241 = arith.xori %ret_205, %ret_240 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_242 = arith.xori %left_idx_229, %right_idx_230 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_243 = arith.select %cond_238, %new_idxs_242, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_244 = arith.xori %new_idxs_208, %new_idxs_243 : tensor<32x16xi32> loc(#loc199)
+    %y_245 = tt.reshape %ret_241 : tensor<32x16xi32> -> tensor<32x2x8xi32> loc(#loc162)
+    %ileft_246 = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<32x2x8xi32> loc(#loc164)
+    %ileft_247 = arith.muli %y_245, %ileft_246 : tensor<32x2x8xi32> loc(#loc164)
+    %ileft_248 = "tt.reduce"(%ileft_247) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc201)
+    %ileft_249 = tt.expand_dims %ileft_248 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc166)
+    %ileft_250 = tt.broadcast %ileft_249 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc167)
+    %iright_251 = arith.muli %y_245, %flip_134 : tensor<32x2x8xi32> loc(#loc168)
+    %iright_252 = "tt.reduce"(%iright_251) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc203)
+    %iright_253 = tt.expand_dims %iright_252 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc170)
+    %iright_254 = tt.broadcast %iright_253 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc171)
+    %ileft_255 = tt.reshape %ileft_250 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_256 = tt.reshape %iright_254 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_257 = tt.reshape %new_idxs_244 : tensor<32x16xi32> -> tensor<32x2x8xi32> loc(#loc174)
+    %left_idx_258 = arith.muli %y_idx_257, %ileft_246 : tensor<32x2x8xi32> loc(#loc176)
+    %left_idx_259 = "tt.reduce"(%left_idx_258) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc206)
+    %left_idx_260 = tt.expand_dims %left_idx_259 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc178)
+    %left_idx_261 = tt.broadcast %left_idx_260 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc179)
+    %right_idx_262 = arith.muli %y_idx_257, %flip_134 : tensor<32x2x8xi32> loc(#loc181)
+    %right_idx_263 = "tt.reduce"(%right_idx_262) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc209)
+    %right_idx_264 = tt.expand_dims %right_idx_263 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc183)
+    %right_idx_265 = tt.broadcast %right_idx_264 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc184)
+    %left_idx_266 = tt.reshape %left_idx_261 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_267 = tt.reshape %right_idx_265 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_268 = arith.cmpi slt, %ileft_255, %iright_256 : tensor<32x16xi32> loc(#loc187)
+    %eq_269 = arith.cmpi eq, %ileft_255, %iright_256 : tensor<32x16xi32> loc(#loc188)
+    %cond_270 = arith.cmpi sgt, %left_idx_266, %right_idx_267 : tensor<32x16xi32> loc(#loc189)
+    %cond_271 = arith.andi %eq_269, %cond_270 : tensor<32x16xi1> loc(#loc190)
+    %cond_272 = arith.ori %cond_268, %cond_271 : tensor<32x16xi1> loc(#loc191)
+    %ret_273 = arith.xori %ileft_255, %iright_256 : tensor<32x16xi32> loc(#loc194)
+    %ret_274 = arith.select %cond_272, %ret_273, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_275 = arith.xori %ret_241, %ret_274 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_276 = arith.xori %left_idx_266, %right_idx_267 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_277 = arith.select %cond_272, %new_idxs_276, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_278 = arith.xori %new_idxs_244, %new_idxs_277 : tensor<32x16xi32> loc(#loc199)
+    %y_279 = tt.reshape %ret_275 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc162)
+    %ileft_280 = arith.muli %y_279, %ileft_137 : tensor<64x2x4xi32> loc(#loc164)
+    %ileft_281 = "tt.reduce"(%ileft_280) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc201)
+    %ileft_282 = tt.expand_dims %ileft_281 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc166)
+    %ileft_283 = tt.broadcast %ileft_282 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc167)
+    %iright_284 = arith.muli %y_279, %flip_59 : tensor<64x2x4xi32> loc(#loc168)
+    %iright_285 = "tt.reduce"(%iright_284) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc203)
+    %iright_286 = tt.expand_dims %iright_285 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc170)
+    %iright_287 = tt.broadcast %iright_286 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc171)
+    %ileft_288 = tt.reshape %ileft_283 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_289 = tt.reshape %iright_287 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_290 = tt.reshape %new_idxs_278 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc174)
+    %left_idx_291 = arith.muli %y_idx_290, %ileft_137 : tensor<64x2x4xi32> loc(#loc176)
+    %left_idx_292 = "tt.reduce"(%left_idx_291) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc206)
+    %left_idx_293 = tt.expand_dims %left_idx_292 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc178)
+    %left_idx_294 = tt.broadcast %left_idx_293 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc179)
+    %right_idx_295 = arith.muli %y_idx_290, %flip_59 : tensor<64x2x4xi32> loc(#loc181)
+    %right_idx_296 = "tt.reduce"(%right_idx_295) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc209)
+    %right_idx_297 = tt.expand_dims %right_idx_296 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc183)
+    %right_idx_298 = tt.broadcast %right_idx_297 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc184)
+    %left_idx_299 = tt.reshape %left_idx_294 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_300 = tt.reshape %right_idx_298 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_301 = arith.cmpi slt, %ileft_288, %iright_289 : tensor<32x16xi32> loc(#loc187)
+    %eq_302 = arith.cmpi eq, %ileft_288, %iright_289 : tensor<32x16xi32> loc(#loc188)
+    %cond_303 = arith.cmpi sgt, %left_idx_299, %right_idx_300 : tensor<32x16xi32> loc(#loc189)
+    %cond_304 = arith.andi %eq_302, %cond_303 : tensor<32x16xi1> loc(#loc190)
+    %cond_305 = arith.ori %cond_301, %cond_304 : tensor<32x16xi1> loc(#loc191)
+    %ret_306 = arith.xori %ileft_288, %iright_289 : tensor<32x16xi32> loc(#loc194)
+    %ret_307 = arith.select %cond_305, %ret_306, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_308 = arith.xori %ret_275, %ret_307 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_309 = arith.xori %left_idx_299, %right_idx_300 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_310 = arith.select %cond_305, %new_idxs_309, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_311 = arith.xori %new_idxs_278, %new_idxs_310 : tensor<32x16xi32> loc(#loc199)
+    %y_312 = tt.reshape %ret_308 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc162)
+    %ileft_313 = arith.muli %y_312, %ileft_62 : tensor<128x2x2xi32> loc(#loc164)
+    %ileft_314 = "tt.reduce"(%ileft_313) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc201)
+    %ileft_315 = tt.expand_dims %ileft_314 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc166)
+    %ileft_316 = tt.broadcast %ileft_315 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc167)
+    %iright_317 = arith.muli %y_312, %flip_22 : tensor<128x2x2xi32> loc(#loc168)
+    %iright_318 = "tt.reduce"(%iright_317) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc203)
+    %iright_319 = tt.expand_dims %iright_318 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc170)
+    %iright_320 = tt.broadcast %iright_319 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc171)
+    %ileft_321 = tt.reshape %ileft_316 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_322 = tt.reshape %iright_320 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_323 = tt.reshape %new_idxs_311 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc174)
+    %left_idx_324 = arith.muli %y_idx_323, %ileft_62 : tensor<128x2x2xi32> loc(#loc176)
+    %left_idx_325 = "tt.reduce"(%left_idx_324) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc206)
+    %left_idx_326 = tt.expand_dims %left_idx_325 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc178)
+    %left_idx_327 = tt.broadcast %left_idx_326 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc179)
+    %right_idx_328 = arith.muli %y_idx_323, %flip_22 : tensor<128x2x2xi32> loc(#loc181)
+    %right_idx_329 = "tt.reduce"(%right_idx_328) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc209)
+    %right_idx_330 = tt.expand_dims %right_idx_329 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc183)
+    %right_idx_331 = tt.broadcast %right_idx_330 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc184)
+    %left_idx_332 = tt.reshape %left_idx_327 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_333 = tt.reshape %right_idx_331 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_334 = arith.cmpi slt, %ileft_321, %iright_322 : tensor<32x16xi32> loc(#loc187)
+    %eq_335 = arith.cmpi eq, %ileft_321, %iright_322 : tensor<32x16xi32> loc(#loc188)
+    %cond_336 = arith.cmpi sgt, %left_idx_332, %right_idx_333 : tensor<32x16xi32> loc(#loc189)
+    %cond_337 = arith.andi %eq_335, %cond_336 : tensor<32x16xi1> loc(#loc190)
+    %cond_338 = arith.ori %cond_334, %cond_337 : tensor<32x16xi1> loc(#loc191)
+    %ret_339 = arith.xori %ileft_321, %iright_322 : tensor<32x16xi32> loc(#loc194)
+    %ret_340 = arith.select %cond_338, %ret_339, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
+    %ret_341 = arith.xori %ret_308, %ret_340 : tensor<32x16xi32> loc(#loc196)
+    %new_idxs_342 = arith.xori %left_idx_332, %right_idx_333 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_343 = arith.select %cond_338, %new_idxs_342, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_344 = arith.xori %new_idxs_311, %new_idxs_343 : tensor<32x16xi32> loc(#loc199)
+    %y_345 = tt.reshape %ret_341 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
+    %ileft_346 = arith.muli %y_345, %ileft : tensor<256x2x1xi32> loc(#loc164)
+    %ileft_347 = "tt.reduce"(%ileft_346) <{axis = 1 : i32}> ({
+    ^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
+      %ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
+      tt.reduce.return %ileft_379 : i32 loc(#loc201)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
+    %ileft_348 = tt.expand_dims %ileft_347 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
+    %ileft_349 = tt.broadcast %ileft_348 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
+    %iright_350 = arith.muli %y_345, %iright : tensor<256x2x1xi32> loc(#loc168)
+    %iright_351 = "tt.reduce"(%iright_350) <{axis = 1 : i32}> ({
+    ^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
+      %iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
+      tt.reduce.return %iright_379 : i32 loc(#loc203)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
+    %iright_352 = tt.expand_dims %iright_351 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
+    %iright_353 = tt.broadcast %iright_352 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
+    %ileft_354 = tt.reshape %ileft_349 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
+    %iright_355 = tt.reshape %iright_353 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
+    %y_idx_356 = tt.reshape %new_idxs_344 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc174)
+    %left_idx_357 = arith.muli %y_idx_356, %ileft : tensor<256x2x1xi32> loc(#loc176)
+    %left_idx_358 = "tt.reduce"(%left_idx_357) <{axis = 1 : i32}> ({
+    ^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
+      %left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
+      tt.reduce.return %left_idx_379 : i32 loc(#loc206)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
+    %left_idx_359 = tt.expand_dims %left_idx_358 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
+    %left_idx_360 = tt.broadcast %left_idx_359 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
+    %right_idx_361 = arith.muli %y_idx_356, %iright : tensor<256x2x1xi32> loc(#loc181)
+    %right_idx_362 = "tt.reduce"(%right_idx_361) <{axis = 1 : i32}> ({
+    ^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
+      %right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
+      tt.reduce.return %right_idx_379 : i32 loc(#loc209)
+    }) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
+    %right_idx_363 = tt.expand_dims %right_idx_362 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
+    %right_idx_364 = tt.broadcast %right_idx_363 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
+    %left_idx_365 = tt.reshape %left_idx_360 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
+    %right_idx_366 = tt.reshape %right_idx_364 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
+    %cond_367 = arith.cmpi slt, %ileft_354, %iright_355 : tensor<32x16xi32> loc(#loc187)
+    %eq_368 = arith.cmpi eq, %ileft_354, %iright_355 : tensor<32x16xi32> loc(#loc188)
+    %cond_369 = arith.cmpi sgt, %left_idx_365, %right_idx_366 : tensor<32x16xi32> loc(#loc189)
+    %cond_370 = arith.andi %eq_368, %cond_369 : tensor<32x16xi1> loc(#loc190)
+    %cond_371 = arith.ori %cond_367, %cond_370 : tensor<32x16xi1> loc(#loc191)
+    %new_idxs_372 = arith.xori %left_idx_365, %right_idx_366 : tensor<32x16xi32> loc(#loc197)
+    %new_idxs_373 = arith.select %cond_371, %new_idxs_372, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
+    %new_idxs_374 = arith.xori %new_idxs_344, %new_idxs_373 : tensor<32x16xi32> loc(#loc199)
+    %tmp7 = arith.extsi %tmp0_19 : tensor<32x16xi32> to tensor<32x16xi64> loc(#loc149)
+    %tmp10_375 = arith.select %tmp0_18, %tmp7, %tmp10 : tensor<32x16xi1>, tensor<32x16xi64> loc(#loc86)
+    %tmp11 = "tt.reduce"(%tmp10_375) <{axis = 1 : i32}> ({
+    ^bb0(%tmp11_377: i64 loc(callsite(#loc1 at #loc150)), %tmp11_378: i64 loc(callsite(#loc1 at #loc150))):
+      %tmp11_379 = arith.addi %tmp11_377, %tmp11_378 : i64 loc(#loc200)
+      tt.reduce.return %tmp11_379 : i64 loc(#loc160)
+    }) : (tensor<32x16xi64>) -> tensor<32xi64> loc(#loc160)
+    %tmp11_376 = tt.expand_dims %tmp11 {axis = 1 : i32} : tensor<32xi64> -> tensor<32x1xi64> loc(#loc151)
+    %tmp14 = arith.trunci %tmp11_376 : tensor<32x1xi64> to tensor<32x1xi32> loc(#loc152)
+    %0 = arith.muli %xindex_6, %cst_2 : tensor<32x1xi32> loc(#loc73)
+    %1 = tt.broadcast %r0_index_8 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc74)
+    %2 = tt.broadcast %0 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc74)
+    %3 = arith.addi %1, %2 : tensor<32x16xi32> loc(#loc74)
+    %4 = tt.splat %out_ptr2 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>> loc(#loc75)
+    %5 = tt.addptr %4, %3 : tensor<32x16x!tt.ptr<i32>>, tensor<32x16xi32> loc(#loc75)
+    tt.store %5, %new_idxs_374, %tmp0_18 : tensor<32x16x!tt.ptr<i32>> loc(#loc76)
+    %6 = tt.splat %out_ptr3 : !tt.ptr<i32> -> tensor<32x1x!tt.ptr<i32>> loc(#loc77)
+    %7 = tt.addptr %6, %xindex_6 : tensor<32x1x!tt.ptr<i32>>, tensor<32x1xi32> loc(#loc77)
+    tt.store %7, %tmp14, %xmask_7 : tensor<32x1x!tt.ptr<i32>> loc(#loc78)
+    tt.return loc(#loc79)
+  } loc(#loc)
+} loc(#loc)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":44:34)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:49)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:38)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":26:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":24:28)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":24:33)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":25:36)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":25:44)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":25:23)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":27:28)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":27:38)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":33:19)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":34:19)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:35)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:45)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:30)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":36:54)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":38:19)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":40:33)
+#loc22 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:41)
+#loc24 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:44)
+#loc25 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:60)
+#loc26 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:68)
+#loc27 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":533:22)
+#loc29 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":537:21)
+#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:40)
+#loc31 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc33 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:65)
+#loc35 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:78)
+#loc36 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:41)
+#loc38 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:67)
+#loc39 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:80)
+#loc40 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":540:30)
+#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":541:32)
+#loc42 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":546:29)
+#loc43 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:36)
+#loc44 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:23)
+#loc45 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":290:25)
+#loc47 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:53)
+#loc48 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:66)
+#loc49 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:37)
+#loc50 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:23)
+#loc52 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:54)
+#loc53 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:67)
+#loc54 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":553:36)
+#loc55 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":554:38)
+#loc56 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":574:22)
+#loc57 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":591:21)
+#loc58 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:40)
+#loc59 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:29)
+#loc60 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:23)
+#loc61 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":599:19)
+#loc62 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":599:28)
+#loc63 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:38)
+#loc64 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:46)
+#loc65 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:15)
+#loc66 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:48)
+#loc67 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:59)
+#loc68 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:22)
+#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":42:19)
+#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":45:29)
+#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":48:21)
+#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:35)
+#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:32)
+#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:25)
+#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":49:47)
+#loc77 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":50:25)
+#loc78 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":50:37)
+#loc79 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/4y/c4yxgotihoxpn6o5xa4jvkcy7shlgnyv44u6dpm5e746f6dwg7oe.py":50:4)
+#loc85 = loc(callsite(#loc1 at #loc2))
+#loc86 = loc("tmp10"(#loc3))
+#loc87 = loc("tmp0"(#loc4))
+#loc88 = loc("tmp0"(#loc5))
+#loc89 = loc("xmask"(#loc6))
+#loc90 = loc("xoffset"(#loc7))
+#loc91 = loc("xoffset"(#loc8))
+#loc92 = loc("xindex"(#loc9))
+#loc93 = loc("xindex"(#loc10))
+#loc94 = loc("xindex"(#loc11))
+#loc95 = loc("r0_index"(#loc12))
+#loc96 = loc("r0_index"(#loc13))
+#loc97 = loc("x0"(#loc14))
+#loc98 = loc("x1"(#loc15))
+#loc99 = loc("tmp0"(#loc16))
+#loc100 = loc("tmp0"(#loc17))
+#loc101 = loc("tmp0"(#loc18))
+#loc102 = loc("tmp0"(#loc19))
+#loc103 = loc("tmp2"(#loc20))
+#loc104 = loc("tmp4"(#loc21))
+#loc105 = loc("flip"(#loc22))
+#loc107 = loc("flip"(#loc24))
+#loc108 = loc("flip"(#loc25))
+#loc109 = loc("flip"(#loc26))
+#loc110 = loc("y"(#loc27))
+#loc111 = loc("left_mask"(#loc29))
+#loc112 = loc("ileft"(#loc30))
+#loc114 = loc("ileft"(#loc34))
+#loc115 = loc("ileft"(#loc35))
+#loc116 = loc("iright"(#loc36))
+#loc118 = loc("iright"(#loc38))
+#loc119 = loc("iright"(#loc39))
+#loc120 = loc("ileft"(#loc40))
+#loc121 = loc("iright"(#loc41))
+#loc122 = loc("y_idx"(#loc42))
+#loc123 = loc("left_idx"(#loc43))
+#loc124 = loc("left_idx"(#loc44))
+#loc125 = loc("input"(#loc45))
+#loc127 = loc("left_idx"(#loc47))
+#loc128 = loc("left_idx"(#loc48))
+#loc129 = loc("right_idx"(#loc49))
+#loc130 = loc("right_idx"(#loc50))
+#loc132 = loc("right_idx"(#loc52))
+#loc133 = loc("right_idx"(#loc53))
+#loc134 = loc("left_idx"(#loc54))
+#loc135 = loc("right_idx"(#loc55))
+#loc136 = loc("cond"(#loc56))
+#loc137 = loc("eq"(#loc57))
+#loc138 = loc("cond"(#loc58))
+#loc139 = loc("cond"(#loc59))
+#loc140 = loc("cond"(#loc60))
+#loc141 = loc("cond"(#loc61))
+#loc142 = loc("cond"(#loc62))
+#loc143 = loc("ret"(#loc63))
+#loc144 = loc("ret"(#loc64))
+#loc145 = loc("ret"(#loc65))
+#loc146 = loc("new_idxs"(#loc66))
+#loc147 = loc("new_idxs"(#loc67))
+#loc148 = loc("new_idxs"(#loc68))
+#loc149 = loc("tmp7"(#loc69))
+#loc151 = loc("tmp11"(#loc71))
+#loc152 = loc("tmp14"(#loc72))
+#loc153 = loc(callsite(#loc105 at #loc106))
+#loc154 = loc(callsite(#loc107 at #loc106))
+#loc155 = loc(callsite(#loc108 at #loc106))
+#loc156 = loc(callsite(#loc109 at #loc106))
+#loc158 = loc("cond"(#loc136))
+#loc159 = loc("eq"(#loc137))
+#loc160 = loc(callsite(#loc31 at #loc150))
+#loc162 = loc(callsite(#loc110 at #loc157))
+#loc163 = loc(callsite(#loc111 at #loc157))
+#loc164 = loc(callsite(#loc112 at #loc157))
+#loc166 = loc(callsite(#loc114 at #loc157))
+#loc167 = loc(callsite(#loc115 at #loc157))
+#loc168 = loc(callsite(#loc116 at #loc157))
+#loc170 = loc(callsite(#loc118 at #loc157))
+#loc171 = loc(callsite(#loc119 at #loc157))
+#loc172 = loc(callsite(#loc120 at #loc157))
+#loc173 = loc(callsite(#loc121 at #loc157))
+#loc174 = loc(callsite(#loc122 at #loc157))
+#loc175 = loc(callsite(#loc123 at #loc157))
+#loc176 = loc(callsite(#loc124 at #loc157))
+#loc178 = loc(callsite(#loc127 at #loc157))
+#loc179 = loc(callsite(#loc128 at #loc157))
+#loc180 = loc(callsite(#loc129 at #loc157))
+#loc181 = loc(callsite(#loc130 at #loc157))
+#loc183 = loc(callsite(#loc132 at #loc157))
+#loc184 = loc(callsite(#loc133 at #loc157))
+#loc185 = loc(callsite(#loc134 at #loc157))
+#loc186 = loc(callsite(#loc135 at #loc157))
+#loc187 = loc(callsite(#loc158 at #loc157))
+#loc188 = loc(callsite(#loc159 at #loc157))
+#loc189 = loc(callsite(#loc138 at #loc157))
+#loc190 = loc(callsite(#loc139 at #loc157))
+#loc191 = loc(callsite(#loc140 at #loc157))
+#loc192 = loc(callsite(#loc141 at #loc157))
+#loc193 = loc(callsite(#loc142 at #loc157))
+#loc194 = loc(callsite(#loc143 at #loc157))
+#loc195 = loc(callsite(#loc144 at #loc157))
+#loc196 = loc(callsite(#loc145 at #loc157))
+#loc197 = loc(callsite(#loc146 at #loc157))
+#loc198 = loc(callsite(#loc147 at #loc157))
+#loc199 = loc(callsite(#loc148 at #loc157))
+#loc200 = loc(callsite(#loc33 at #loc160))
+#loc201 = loc(callsite(#loc31 at #loc165))
+#loc203 = loc(callsite(#loc31 at #loc169))
+#loc205 = loc(callsite(#loc125 at #loc177))
+#loc206 = loc(callsite(#loc31 at #loc177))
+#loc208 = loc(callsite(#loc125 at #loc182))
+#loc209 = loc(callsite(#loc31 at #loc182))
+#loc211 = loc(callsite(#loc33 at #loc201))
+#loc212 = loc(callsite(#loc33 at #loc203))
+#loc213 = loc(callsite(#loc33 at #loc206))
+#loc214 = loc(callsite(#loc33 at #loc209))

SpecForge-ext/cache/compiled_kernels/triton/6/7AX37JF3GC3476WMWTYVB3NGNRIS5J7SJEWETBMPGJB4H67EAVCA/triton_per_fused__to_copy_mul_sum_1.llir

@@ -0,0 +1,220 @@
+; ModuleID = 'LLVMDialectModule'
+source_filename = "LLVMDialectModule"
+target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+
+@global_smem = external addrspace(3) global [0 x i8], align 16
+
+; Function Attrs: nounwind
+define ptx_kernel void @triton_per_fused__to_copy_mul_sum_1(ptr addrspace(1) %0, ptr addrspace(1) %1, i64 %2, i32 %3, i32 %4, ptr addrspace(1) readnone captures(none) %5, ptr addrspace(1) readnone captures(none) %6) local_unnamed_addr #0 !dbg !4 {
+  %8 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !7
+  %9 = shl i32 %8, 7, !dbg !8
+  %10 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !9
+  %11 = and i32 %10, 127, !dbg !9
+  %12 = or disjoint i32 %9, %11, !dbg !10
+  %13 = icmp slt i32 %12, %3, !dbg !11
+  %14 = lshr i32 %10, 7, !dbg !12
+  %.lobit = and i32 %14, 1, !dbg !12
+  %15 = or disjoint i32 %.lobit, 2, !dbg !12
+  %16 = or disjoint i32 %.lobit, 4, !dbg !12
+  %17 = or i32 %14, 6, !dbg !12
+  %18 = or disjoint i32 %.lobit, 8, !dbg !12
+  %19 = or disjoint i32 %.lobit, 10, !dbg !12
+  %20 = or disjoint i32 %.lobit, 12, !dbg !12
+  %21 = or i32 %14, 14, !dbg !12
+  %22 = or disjoint i32 %.lobit, 16, !dbg !12
+  %23 = or disjoint i32 %.lobit, 18, !dbg !12
+  %24 = or disjoint i32 %.lobit, 20, !dbg !12
+  %25 = or i32 %14, 22, !dbg !12
+  %26 = or disjoint i32 %.lobit, 24, !dbg !12
+  %27 = or disjoint i32 %.lobit, 26, !dbg !12
+  %28 = or disjoint i32 %.lobit, 28, !dbg !12
+  %29 = or i32 %14, 30, !dbg !12
+  %30 = zext nneg i32 %.lobit to i64, !dbg !13
+  %31 = zext nneg i32 %15 to i64, !dbg !13
+  %32 = zext nneg i32 %16 to i64, !dbg !13
+  %33 = zext nneg i32 %17 to i64, !dbg !13
+  %34 = zext nneg i32 %18 to i64, !dbg !13
+  %35 = zext nneg i32 %19 to i64, !dbg !13
+  %36 = zext nneg i32 %20 to i64, !dbg !13
+  %37 = zext nneg i32 %21 to i64, !dbg !13
+  %38 = zext nneg i32 %22 to i64, !dbg !13
+  %39 = zext nneg i32 %23 to i64, !dbg !13
+  %40 = zext nneg i32 %24 to i64, !dbg !13
+  %41 = zext nneg i32 %25 to i64, !dbg !13
+  %42 = zext nneg i32 %26 to i64, !dbg !13
+  %43 = zext nneg i32 %27 to i64, !dbg !13
+  %44 = zext nneg i32 %28 to i64, !dbg !13
+  %45 = zext nneg i32 %29 to i64, !dbg !13
+  %46 = mul nuw nsw i64 %2, %30, !dbg !13
+  %47 = mul i64 %2, %31, !dbg !13
+  %48 = mul i64 %2, %32, !dbg !13
+  %49 = mul i64 %2, %33, !dbg !13
+  %50 = mul i64 %2, %34, !dbg !13
+  %51 = mul i64 %2, %35, !dbg !13
+  %52 = mul i64 %2, %36, !dbg !13
+  %53 = mul i64 %2, %37, !dbg !13
+  %54 = mul i64 %2, %38, !dbg !13
+  %55 = mul i64 %2, %39, !dbg !13
+  %56 = mul i64 %2, %40, !dbg !13
+  %57 = mul i64 %2, %41, !dbg !13
+  %58 = mul i64 %2, %42, !dbg !13
+  %59 = mul i64 %2, %43, !dbg !13
+  %60 = mul i64 %2, %44, !dbg !13
+  %61 = mul i64 %2, %45, !dbg !13
+  %62 = sext i32 %12 to i64, !dbg !14
+  %63 = getelementptr float, ptr addrspace(1) %0, i64 %46, !dbg !15
+  %64 = getelementptr float, ptr addrspace(1) %63, i64 %62, !dbg !15
+  %65 = getelementptr float, ptr addrspace(1) %0, i64 %47, !dbg !15
+  %66 = getelementptr float, ptr addrspace(1) %65, i64 %62, !dbg !15
+  %67 = getelementptr float, ptr addrspace(1) %0, i64 %48, !dbg !15
+  %68 = getelementptr float, ptr addrspace(1) %67, i64 %62, !dbg !15
+  %69 = getelementptr float, ptr addrspace(1) %0, i64 %49, !dbg !15
+  %70 = getelementptr float, ptr addrspace(1) %69, i64 %62, !dbg !15
+  %71 = getelementptr float, ptr addrspace(1) %0, i64 %50, !dbg !15
+  %72 = getelementptr float, ptr addrspace(1) %71, i64 %62, !dbg !15
+  %73 = getelementptr float, ptr addrspace(1) %0, i64 %51, !dbg !15
+  %74 = getelementptr float, ptr addrspace(1) %73, i64 %62, !dbg !15
+  %75 = getelementptr float, ptr addrspace(1) %0, i64 %52, !dbg !15
+  %76 = getelementptr float, ptr addrspace(1) %75, i64 %62, !dbg !15
+  %77 = getelementptr float, ptr addrspace(1) %0, i64 %53, !dbg !15
+  %78 = getelementptr float, ptr addrspace(1) %77, i64 %62, !dbg !15
+  %79 = getelementptr float, ptr addrspace(1) %0, i64 %54, !dbg !15
+  %80 = getelementptr float, ptr addrspace(1) %79, i64 %62, !dbg !15
+  %81 = getelementptr float, ptr addrspace(1) %0, i64 %55, !dbg !15
+  %82 = getelementptr float, ptr addrspace(1) %81, i64 %62, !dbg !15
+  %83 = getelementptr float, ptr addrspace(1) %0, i64 %56, !dbg !15
+  %84 = getelementptr float, ptr addrspace(1) %83, i64 %62, !dbg !15
+  %85 = getelementptr float, ptr addrspace(1) %0, i64 %57, !dbg !15
+  %86 = getelementptr float, ptr addrspace(1) %85, i64 %62, !dbg !15
+  %87 = getelementptr float, ptr addrspace(1) %0, i64 %58, !dbg !15
+  %88 = getelementptr float, ptr addrspace(1) %87, i64 %62, !dbg !15
+  %89 = getelementptr float, ptr addrspace(1) %0, i64 %59, !dbg !15
+  %90 = getelementptr float, ptr addrspace(1) %89, i64 %62, !dbg !15
+  %91 = getelementptr float, ptr addrspace(1) %0, i64 %60, !dbg !15
+  %92 = getelementptr float, ptr addrspace(1) %91, i64 %62, !dbg !15
+  %93 = getelementptr float, ptr addrspace(1) %0, i64 %61, !dbg !15
+  %94 = getelementptr float, ptr addrspace(1) %93, i64 %62, !dbg !15
+  %95 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %64, i1 %13) #4, !dbg !16
+  %96 = bitcast i32 %95 to float, !dbg !16
+  %97 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %66, i1 %13) #4, !dbg !16
+  %98 = bitcast i32 %97 to float, !dbg !16
+  %99 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %68, i1 %13) #4, !dbg !16
+  %100 = bitcast i32 %99 to float, !dbg !16
+  %101 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %70, i1 %13) #4, !dbg !16
+  %102 = bitcast i32 %101 to float, !dbg !16
+  %103 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %72, i1 %13) #4, !dbg !16
+  %104 = bitcast i32 %103 to float, !dbg !16
+  %105 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %74, i1 %13) #4, !dbg !16
+  %106 = bitcast i32 %105 to float, !dbg !16
+  %107 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %76, i1 %13) #4, !dbg !16
+  %108 = bitcast i32 %107 to float, !dbg !16
+  %109 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %78, i1 %13) #4, !dbg !16
+  %110 = bitcast i32 %109 to float, !dbg !16
+  %111 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %80, i1 %13) #4, !dbg !16
+  %112 = bitcast i32 %111 to float, !dbg !16
+  %113 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %82, i1 %13) #4, !dbg !16
+  %114 = bitcast i32 %113 to float, !dbg !16
+  %115 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %84, i1 %13) #4, !dbg !16
+  %116 = bitcast i32 %115 to float, !dbg !16
+  %117 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %86, i1 %13) #4, !dbg !16
+  %118 = bitcast i32 %117 to float, !dbg !16
+  %119 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %88, i1 %13) #4, !dbg !16
+  %120 = bitcast i32 %119 to float, !dbg !16
+  %121 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %90, i1 %13) #4, !dbg !16
+  %122 = bitcast i32 %121 to float, !dbg !16
+  %123 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %92, i1 %13) #4, !dbg !16
+  %124 = bitcast i32 %123 to float, !dbg !16
+  %125 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %94, i1 %13) #4, !dbg !16
+  %126 = bitcast i32 %125 to float, !dbg !16
+  %127 = fadd float %96, %98, !dbg !17
+  %128 = fadd float %127, %100, !dbg !17
+  %129 = fadd float %128, %102, !dbg !17
+  %130 = fadd float %129, %104, !dbg !17
+  %131 = fadd float %130, %106, !dbg !17
+  %132 = fadd float %131, %108, !dbg !17
+  %133 = fadd float %132, %110, !dbg !17
+  %134 = fadd float %133, %112, !dbg !17
+  %135 = fadd float %134, %114, !dbg !17
+  %136 = fadd float %135, %116, !dbg !17
+  %137 = fadd float %136, %118, !dbg !17
+  %138 = fadd float %137, %120, !dbg !17
+  %139 = fadd float %138, %122, !dbg !17
+  %140 = fadd float %139, %124, !dbg !17
+  %141 = fadd float %140, %126, !dbg !17
+  %.idx = shl nuw nsw i32 %11, 3, !dbg !21
+  %142 = getelementptr i8, ptr addrspace(3) @global_smem, i32 %.idx, !dbg !21
+  %143 = getelementptr float, ptr addrspace(3) %142, i32 %.lobit, !dbg !21
+  %144 = bitcast float %141 to i32, !dbg !21
+  %145 = select i1 %13, i32 %144, i32 0, !dbg !17
+  %146 = insertelement <1 x i32> poison, i32 %145, i64 0, !dbg !21
+  tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %143, <1 x i32> %146, i1 true) #4, !dbg !21
+  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !21
+  %147 = icmp samesign ult i32 %10, 256, !dbg !21
+  %148 = getelementptr float, ptr addrspace(3) @global_smem, i32 %10, !dbg !21
+  %149 = tail call i32 asm sideeffect "@$2 ld.shared.b32 $0, [ $1 + 0 ];", "=r,r,b"(ptr addrspace(3) %148, i1 %147) #4, !dbg !21
+  %150 = bitcast i32 %149 to float, !dbg !21
+  %151 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %149, i32 1, i32 31), !dbg !21
+  %152 = bitcast i32 %151 to float, !dbg !21
+  %153 = fadd float %150, %152, !dbg !17
+  %154 = and i32 %10, 769, !dbg !21
+  %155 = icmp eq i32 %154, 0, !dbg !21
+  %156 = bitcast float %153 to <1 x i32>, !dbg !21
+  tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %148, <1 x i32> %156, i1 %155) #4, !dbg !21
+  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !21
+  %157 = load float, ptr addrspace(3) %142, align 8, !dbg !21
+  %158 = getelementptr bfloat, ptr addrspace(1) %1, i64 %62, !dbg !22
+  %159 = fptrunc float %157 to bfloat, !dbg !23
+  %160 = and i32 %10, 128, !dbg !23
+  %161 = icmp eq i32 %160, 0, !dbg !23
+  %162 = bitcast bfloat %159 to i16, !dbg !23
+  %163 = and i1 %161, %13, !dbg !23
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %162, ptr addrspace(1) %158, i1 %163) #4, !dbg !23
+  ret void, !dbg !24
+}
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
+
+; Function Attrs: convergent nocallback nounwind
+declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #2
+
+; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
+declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) #3
+
+attributes #0 = { nounwind "nvvm.reqntid"="256" }
+attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+attributes #2 = { convergent nocallback nounwind }
+attributes #3 = { convergent nocallback nounwind memory(inaccessiblemem: readwrite) }
+attributes #4 = { nounwind }
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
+!1 = !DIFile(filename: "cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf")
+!2 = !{i32 2, !"Debug Info Version", i32 3}
+!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
+!4 = distinct !DISubprogram(name: "triton_per_fused__to_copy_mul_sum_1", linkageName: "triton_per_fused__to_copy_mul_sum_1", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
+!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
+!6 = !{}
+!7 = !DILocation(line: 23, column: 28, scope: !4)
+!8 = !DILocation(line: 23, column: 33, scope: !4)
+!9 = !DILocation(line: 24, column: 44, scope: !4)
+!10 = !DILocation(line: 24, column: 23, scope: !4)
+!11 = !DILocation(line: 25, column: 21, scope: !4)
+!12 = !DILocation(line: 26, column: 38, scope: !4)
+!13 = !DILocation(line: 33, column: 39, scope: !4)
+!14 = !DILocation(line: 33, column: 35, scope: !4)
+!15 = !DILocation(line: 33, column: 30, scope: !4)
+!16 = !DILocation(line: 33, column: 46, scope: !4)
+!17 = !DILocation(line: 261, column: 15, scope: !18, inlinedAt: !20)
+!18 = distinct !DILexicalBlockFile(scope: !4, file: !19, discriminator: 0)
+!19 = !DIFile(filename: "standard.py", directory: "/workspace/specforge/lib/python3.11/site-packages/triton/language")
+!20 = !DILocation(line: 36, column: 24, scope: !4)
+!21 = !DILocation(line: 291, column: 36, scope: !18, inlinedAt: !20)
+!22 = !DILocation(line: 37, column: 25, scope: !4)
+!23 = !DILocation(line: 37, column: 36, scope: !4)
+!24 = !DILocation(line: 37, column: 4, scope: !4)

SpecForge-ext/cache/compiled_kernels/triton/6/7AX37JF3GC3476WMWTYVB3NGNRIS5J7SJEWETBMPGJB4H67EAVCA/triton_per_fused__to_copy_mul_sum_1.ptx

@@ -0,0 +1,522 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.7
+.target sm_90a
+.address_size 64
+
+	// .globl	triton_per_fused__to_copy_mul_sum_1 // -- Begin function triton_per_fused__to_copy_mul_sum_1
+.extern .shared .align 16 .b8 global_smem[];
+                                        // @triton_per_fused__to_copy_mul_sum_1
+.visible .entry triton_per_fused__to_copy_mul_sum_1(
+	.param .u64 .ptr .global .align 1 triton_per_fused__to_copy_mul_sum_1_param_0,
+	.param .u64 .ptr .global .align 1 triton_per_fused__to_copy_mul_sum_1_param_1,
+	.param .u64 triton_per_fused__to_copy_mul_sum_1_param_2,
+	.param .u32 triton_per_fused__to_copy_mul_sum_1_param_3,
+	.param .u32 triton_per_fused__to_copy_mul_sum_1_param_4,
+	.param .u64 .ptr .global .align 1 triton_per_fused__to_copy_mul_sum_1_param_5,
+	.param .u64 .ptr .global .align 1 triton_per_fused__to_copy_mul_sum_1_param_6
+)
+.reqntid 256
+{
+	.reg .pred 	%p<22>;
+	.reg .b16 	%rs<2>;
+	.reg .b32 	%r<86>;
+	.reg .b64 	%rd<86>;
+	.loc	1 18 0                          // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:18:0
+$L__func_begin0:
+	.loc	1 18 0                          // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:18:0
+
+// %bb.0:
+	ld.param.b64 	%rd18, [triton_per_fused__to_copy_mul_sum_1_param_0];
+	ld.param.b64 	%rd19, [triton_per_fused__to_copy_mul_sum_1_param_1];
+$L__tmp0:
+	.loc	1 23 28                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:23:28
+	mov.u32 	%r39, %ctaid.x;
+	.loc	1 23 33                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:23:33
+	shl.b32 	%r40, %r39, 7;
+	ld.param.b64 	%rd20, [triton_per_fused__to_copy_mul_sum_1_param_2];
+	ld.param.b32 	%r41, [triton_per_fused__to_copy_mul_sum_1_param_3];
+	.loc	1 24 44                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:24:44
+	mov.u32 	%r42, %tid.x;
+	and.b32 	%r43, %r42, 127;
+	.loc	1 24 23                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:24:23
+	or.b32 	%r44, %r40, %r43;
+	.loc	1 25 21                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:25:21
+	setp.lt.s32 	%p1, %r44, %r41;
+	.loc	1 26 38                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:26:38
+	shr.u32 	%r45, %r42, 7;
+	bfe.u32 	%r46, %r42, 7, 1;
+	or.b32 	%r47, %r46, 2;
+	or.b32 	%r48, %r46, 4;
+	or.b32 	%r49, %r45, 6;
+	or.b32 	%r50, %r46, 8;
+	or.b32 	%r51, %r46, 10;
+	or.b32 	%r52, %r46, 12;
+	or.b32 	%r53, %r45, 14;
+	or.b32 	%r54, %r46, 16;
+	or.b32 	%r55, %r46, 18;
+	or.b32 	%r56, %r46, 20;
+	or.b32 	%r57, %r45, 22;
+	or.b32 	%r58, %r46, 24;
+	or.b32 	%r59, %r46, 26;
+	or.b32 	%r60, %r46, 28;
+	or.b32 	%r61, %r45, 30;
+	.loc	1 33 39                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:33:39
+	cvt.u64.u32 	%rd21, %r46;
+	cvt.u64.u32 	%rd22, %r47;
+	cvt.u64.u32 	%rd23, %r48;
+	cvt.u64.u32 	%rd24, %r49;
+	cvt.u64.u32 	%rd25, %r50;
+	cvt.u64.u32 	%rd26, %r51;
+	cvt.u64.u32 	%rd27, %r52;
+	cvt.u64.u32 	%rd28, %r53;
+	cvt.u64.u32 	%rd29, %r54;
+	cvt.u64.u32 	%rd30, %r55;
+	cvt.u64.u32 	%rd31, %r56;
+	cvt.u64.u32 	%rd32, %r57;
+	cvt.u64.u32 	%rd33, %r58;
+	cvt.u64.u32 	%rd34, %r59;
+	cvt.u64.u32 	%rd35, %r60;
+	cvt.u64.u32 	%rd36, %r61;
+	mul.lo.s64 	%rd37, %rd20, %rd21;
+	mul.lo.s64 	%rd38, %rd20, %rd22;
+	mul.lo.s64 	%rd39, %rd20, %rd23;
+	mul.lo.s64 	%rd40, %rd20, %rd24;
+	mul.lo.s64 	%rd41, %rd20, %rd25;
+	mul.lo.s64 	%rd42, %rd20, %rd26;
+	mul.lo.s64 	%rd43, %rd20, %rd27;
+	mul.lo.s64 	%rd44, %rd20, %rd28;
+	mul.lo.s64 	%rd45, %rd20, %rd29;
+	mul.lo.s64 	%rd46, %rd20, %rd30;
+	mul.lo.s64 	%rd47, %rd20, %rd31;
+	mul.lo.s64 	%rd48, %rd20, %rd32;
+	mul.lo.s64 	%rd49, %rd20, %rd33;
+	mul.lo.s64 	%rd50, %rd20, %rd34;
+	mul.lo.s64 	%rd51, %rd20, %rd35;
+	mul.lo.s64 	%rd52, %rd20, %rd36;
+	.loc	1 33 30                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:33:30
+	shl.b64 	%rd53, %rd37, 2;
+	add.s64 	%rd54, %rd18, %rd53;
+	mul.wide.s32 	%rd55, %r44, 4;
+	add.s64 	%rd1, %rd54, %rd55;
+	shl.b64 	%rd56, %rd38, 2;
+	add.s64 	%rd57, %rd18, %rd56;
+	add.s64 	%rd2, %rd57, %rd55;
+	shl.b64 	%rd58, %rd39, 2;
+	add.s64 	%rd59, %rd18, %rd58;
+	add.s64 	%rd3, %rd59, %rd55;
+	shl.b64 	%rd60, %rd40, 2;
+	add.s64 	%rd61, %rd18, %rd60;
+	add.s64 	%rd4, %rd61, %rd55;
+	shl.b64 	%rd62, %rd41, 2;
+	add.s64 	%rd63, %rd18, %rd62;
+	add.s64 	%rd5, %rd63, %rd55;
+	shl.b64 	%rd64, %rd42, 2;
+	add.s64 	%rd65, %rd18, %rd64;
+	add.s64 	%rd6, %rd65, %rd55;
+	shl.b64 	%rd66, %rd43, 2;
+	add.s64 	%rd67, %rd18, %rd66;
+	add.s64 	%rd7, %rd67, %rd55;
+	shl.b64 	%rd68, %rd44, 2;
+	add.s64 	%rd69, %rd18, %rd68;
+	add.s64 	%rd8, %rd69, %rd55;
+	shl.b64 	%rd70, %rd45, 2;
+	add.s64 	%rd71, %rd18, %rd70;
+	add.s64 	%rd9, %rd71, %rd55;
+	shl.b64 	%rd72, %rd46, 2;
+	add.s64 	%rd73, %rd18, %rd72;
+	add.s64 	%rd10, %rd73, %rd55;
+	shl.b64 	%rd74, %rd47, 2;
+	add.s64 	%rd75, %rd18, %rd74;
+	add.s64 	%rd11, %rd75, %rd55;
+	shl.b64 	%rd76, %rd48, 2;
+	add.s64 	%rd77, %rd18, %rd76;
+	add.s64 	%rd12, %rd77, %rd55;
+	shl.b64 	%rd78, %rd49, 2;
+	add.s64 	%rd79, %rd18, %rd78;
+	add.s64 	%rd13, %rd79, %rd55;
+	shl.b64 	%rd80, %rd50, 2;
+	add.s64 	%rd81, %rd18, %rd80;
+	add.s64 	%rd14, %rd81, %rd55;
+	shl.b64 	%rd82, %rd51, 2;
+	add.s64 	%rd83, %rd18, %rd82;
+	add.s64 	%rd15, %rd83, %rd55;
+	shl.b64 	%rd84, %rd52, 2;
+	add.s64 	%rd85, %rd18, %rd84;
+	add.s64 	%rd16, %rd85, %rd55;
+	mov.b32 	%r2, 0;
+	.loc	1 33 46                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:33:46
+	// begin inline asm
+	mov.u32 %r1, %r2;
+	@%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r3, %r2;
+	@%p1 ld.global.b32 { %r3 }, [ %rd2 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r5, %r2;
+	@%p1 ld.global.b32 { %r5 }, [ %rd3 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r7, %r2;
+	@%p1 ld.global.b32 { %r7 }, [ %rd4 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r9, %r2;
+	@%p1 ld.global.b32 { %r9 }, [ %rd5 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r11, %r2;
+	@%p1 ld.global.b32 { %r11 }, [ %rd6 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r13, %r2;
+	@%p1 ld.global.b32 { %r13 }, [ %rd7 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r15, %r2;
+	@%p1 ld.global.b32 { %r15 }, [ %rd8 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r17, %r2;
+	@%p1 ld.global.b32 { %r17 }, [ %rd9 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r19, %r2;
+	@%p1 ld.global.b32 { %r19 }, [ %rd10 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r21, %r2;
+	@%p1 ld.global.b32 { %r21 }, [ %rd11 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r23, %r2;
+	@%p1 ld.global.b32 { %r23 }, [ %rd12 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r25, %r2;
+	@%p1 ld.global.b32 { %r25 }, [ %rd13 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r27, %r2;
+	@%p1 ld.global.b32 { %r27 }, [ %rd14 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r29, %r2;
+	@%p1 ld.global.b32 { %r29 }, [ %rd15 + 0 ];
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r31, %r2;
+	@%p1 ld.global.b32 { %r31 }, [ %rd16 + 0 ];
+	// end inline asm
+$L__tmp1:
+	.loc	2 261 15                        // standard.py:261:15 @[ cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:36:24 ]
+	add.f32 	%r62, %r1, %r3;
+	add.f32 	%r63, %r62, %r5;
+	add.f32 	%r64, %r63, %r7;
+	add.f32 	%r65, %r64, %r9;
+	add.f32 	%r66, %r65, %r11;
+	add.f32 	%r67, %r66, %r13;
+	add.f32 	%r68, %r67, %r15;
+	add.f32 	%r69, %r68, %r17;
+	add.f32 	%r70, %r69, %r19;
+	add.f32 	%r71, %r70, %r21;
+	add.f32 	%r72, %r71, %r23;
+	add.f32 	%r73, %r72, %r25;
+	add.f32 	%r74, %r73, %r27;
+	add.f32 	%r75, %r74, %r29;
+	add.f32 	%r76, %r75, %r31;
+	.loc	2 291 36                        // standard.py:291:36 @[ cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:36:24 ]
+	shl.b32 	%r77, %r43, 3;
+	mov.b32 	%r78, global_smem;
+	add.s32 	%r79, %r78, %r77;
+	shl.b32 	%r80, %r46, 2;
+	add.s32 	%r33, %r79, %r80;
+	.loc	2 261 15                        // standard.py:261:15 @[ cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:36:24 ]
+	selp.b32 	%r34, %r76, 0, %p1;
+	mov.pred 	%p17, -1;
+	.loc	2 291 36                        // standard.py:291:36 @[ cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:36:24 ]
+	// begin inline asm
+	@%p17 st.shared.b32 [ %r33 + 0 ], %r34;
+	// end inline asm
+	bar.sync 	0;
+	setp.lt.u32 	%p18, %r42, 256;
+	shl.b32 	%r81, %r42, 2;
+	add.s32 	%r36, %r78, %r81;
+	// begin inline asm
+	@%p18 ld.shared.b32 %r35, [ %r36 + 0 ];
+	// end inline asm
+	shfl.sync.bfly.b32 	%r82, %r35, 1, 31, -1;
+	.loc	2 261 15                        // standard.py:261:15 @[ cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:36:24 ]
+	add.f32 	%r38, %r35, %r82;
+	.loc	2 291 36                        // standard.py:291:36 @[ cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:36:24 ]
+	and.b32 	%r83, %r42, 769;
+	setp.eq.b32 	%p19, %r83, 0;
+	// begin inline asm
+	@%p19 st.shared.b32 [ %r36 + 0 ], %r38;
+	// end inline asm
+	bar.sync 	0;
+	ld.shared.b32 	%r84, [%r79];
+$L__tmp2:
+	.loc	1 37 25                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:37:25
+	mad.wide.s32 	%rd17, %r44, 2, %rd19;
+	.loc	1 37 36                         // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:37:36
+	cvt.rn.bf16.f32 	%rs1, %r84;
+	and.b32 	%r85, %r42, 128;
+	setp.eq.b32 	%p21, %r85, 0;
+	and.pred 	%p20, %p21, %p1;
+	// begin inline asm
+	@%p20 st.global.b16 [ %rd17 + 0 ], { %rs1 };
+	// end inline asm
+	.loc	1 37 4                          // cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py:37:4
+	ret;
+$L__tmp3:
+$L__func_end0:
+                                        // -- End function
+}
+	.file	1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py"
+	.file	2 "/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py"
+	.section	.debug_abbrev
+	{
+.b8 1                                   // Abbreviation Code
+.b8 17                                  // DW_TAG_compile_unit
+.b8 1                                   // DW_CHILDREN_yes
+.b8 37                                  // DW_AT_producer
+.b8 8                                   // DW_FORM_string
+.b8 19                                  // DW_AT_language
+.b8 5                                   // DW_FORM_data2
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 16                                  // DW_AT_stmt_list
+.b8 6                                   // DW_FORM_data4
+.b8 27                                  // DW_AT_comp_dir
+.b8 8                                   // DW_FORM_string
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 2                                   // Abbreviation Code
+.b8 46                                  // DW_TAG_subprogram
+.b8 0                                   // DW_CHILDREN_no
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 32                                  // DW_AT_inline
+.b8 11                                  // DW_FORM_data1
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 3                                   // Abbreviation Code
+.b8 46                                  // DW_TAG_subprogram
+.b8 1                                   // DW_CHILDREN_yes
+.b8 17                                  // DW_AT_low_pc
+.b8 1                                   // DW_FORM_addr
+.b8 18                                  // DW_AT_high_pc
+.b8 1                                   // DW_FORM_addr
+.b8 49                                  // DW_AT_abstract_origin
+.b8 19                                  // DW_FORM_ref4
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 4                                   // Abbreviation Code
+.b8 29                                  // DW_TAG_inlined_subroutine
+.b8 0                                   // DW_CHILDREN_no
+.b8 49                                  // DW_AT_abstract_origin
+.b8 19                                  // DW_FORM_ref4
+.b8 17                                  // DW_AT_low_pc
+.b8 1                                   // DW_FORM_addr
+.b8 18                                  // DW_AT_high_pc
+.b8 1                                   // DW_FORM_addr
+.b8 88                                  // DW_AT_call_file
+.b8 11                                  // DW_FORM_data1
+.b8 89                                  // DW_AT_call_line
+.b8 11                                  // DW_FORM_data1
+.b8 87                                  // DW_AT_call_column
+.b8 11                                  // DW_FORM_data1
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 0                                   // EOM(3)
+	}
+	.section	.debug_info
+	{
+.b32 220                                // Length of Unit
+.b8 2                                   // DWARF version number
+.b8 0
+.b32 .debug_abbrev                      // Offset Into Abbrev. Section
+.b8 8                                   // Address Size (in bytes)
+.b8 1                                   // Abbrev [1] 0xb:0xd5 DW_TAG_compile_unit
+.b8 116                                 // DW_AT_producer
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2                                   // DW_AT_language
+.b8 0
+.b8 99                                  // DW_AT_name
+.b8 121
+.b8 102
+.b8 116
+.b8 55
+.b8 115
+.b8 105
+.b8 97
+.b8 108
+.b8 101
+.b8 112
+.b8 114
+.b8 105
+.b8 119
+.b8 54
+.b8 101
+.b8 117
+.b8 106
+.b8 117
+.b8 108
+.b8 97
+.b8 120
+.b8 112
+.b8 105
+.b8 53
+.b8 55
+.b8 113
+.b8 108
+.b8 114
+.b8 97
+.b8 102
+.b8 107
+.b8 109
+.b8 112
+.b8 52
+.b8 107
+.b8 50
+.b8 107
+.b8 106
+.b8 119
+.b8 122
+.b8 119
+.b8 52
+.b8 110
+.b8 111
+.b8 104
+.b8 50
+.b8 51
+.b8 100
+.b8 100
+.b8 122
+.b8 54
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line                        // DW_AT_stmt_list
+.b8 47                                  // DW_AT_comp_dir
+.b8 119
+.b8 111
+.b8 114
+.b8 107
+.b8 115
+.b8 112
+.b8 97
+.b8 99
+.b8 101
+.b8 47
+.b8 104
+.b8 97
+.b8 110
+.b8 114
+.b8 117
+.b8 105
+.b8 47
+.b8 83
+.b8 112
+.b8 101
+.b8 99
+.b8 70
+.b8 111
+.b8 114
+.b8 103
+.b8 101
+.b8 45
+.b8 101
+.b8 120
+.b8 116
+.b8 47
+.b8 99
+.b8 97
+.b8 99
+.b8 104
+.b8 101
+.b8 47
+.b8 99
+.b8 111
+.b8 109
+.b8 112
+.b8 105
+.b8 108
+.b8 101
+.b8 100
+.b8 95
+.b8 107
+.b8 101
+.b8 114
+.b8 110
+.b8 101
+.b8 108
+.b8 115
+.b8 47
+.b8 121
+.b8 102
+.b8 0
+.b8 2                                   // Abbrev [2] 0x8b:0x26 DW_TAG_subprogram
+.b8 116                                 // DW_AT_name
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 112
+.b8 101
+.b8 114
+.b8 95
+.b8 102
+.b8 117
+.b8 115
+.b8 101
+.b8 100
+.b8 95
+.b8 95
+.b8 116
+.b8 111
+.b8 95
+.b8 99
+.b8 111
+.b8 112
+.b8 121
+.b8 95
+.b8 109
+.b8 117
+.b8 108
+.b8 95
+.b8 115
+.b8 117
+.b8 109
+.b8 95
+.b8 49
+.b8 0
+.b8 1                                   // DW_AT_inline
+.b8 3                                   // Abbrev [3] 0xb1:0x2e DW_TAG_subprogram
+.b64 $L__func_begin0                    // DW_AT_low_pc
+.b64 $L__func_end0                      // DW_AT_high_pc
+.b32 139                                // DW_AT_abstract_origin
+.b8 4                                   // Abbrev [4] 0xc6:0x18 DW_TAG_inlined_subroutine
+.b32 139                                // DW_AT_abstract_origin
+.b64 $L__tmp1                           // DW_AT_low_pc
+.b64 $L__tmp2                           // DW_AT_high_pc
+.b8 1                                   // DW_AT_call_file
+.b8 36                                  // DW_AT_call_line
+.b8 24                                  // DW_AT_call_column
+.b8 0                                   // End Of Children Mark
+.b8 0                                   // End Of Children Mark
+	}
+	.section	.debug_macinfo	{	}

SpecForge-ext/cache/compiled_kernels/triton/6/7AX37JF3GC3476WMWTYVB3NGNRIS5J7SJEWETBMPGJB4H67EAVCA/triton_per_fused__to_copy_mul_sum_1.source

@@ -0,0 +1,120 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":18:0)
+#loc22 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":285:0)
+#loc24 = loc(unknown)
+#loc27 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":260:0)
+#loc31 = loc("in_ptr0"(#loc))
+#loc32 = loc("out_ptr0"(#loc))
+#loc33 = loc("ks0"(#loc))
+#loc34 = loc("xnumel"(#loc))
+#loc35 = loc("r0_numel"(#loc))
+#loc54 = loc("input"(#loc22))
+#loc55 = loc("a"(#loc27))
+#loc56 = loc("b"(#loc27))
+module {
+  tt.func public @triton_per_fused__to_copy_mul_sum_1(%in_ptr0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %xnumel: i32 loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %r0_numel_0 = arith.constant 32 : i32 loc(#loc36)
+    %xoffset = tt.get_program_id x : i32 loc(#loc37)
+    %xoffset_1 = arith.constant 128 : i32 loc(#loc38)
+    %xoffset_2 = arith.constant 128 : i32 loc(#loc38)
+    %xoffset_3 = arith.muli %xoffset, %xoffset_2 : i32 loc(#loc38)
+    %xindex = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc39)
+    %xindex_4 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<128xi32> -> tensor<128x1xi32> loc(#loc40)
+    %xindex_5 = tt.splat %xoffset_3 : i32 -> tensor<128x1xi32> loc(#loc41)
+    %xindex_6 = arith.addi %xindex_5, %xindex_4 : tensor<128x1xi32> loc(#loc41)
+    %xmask = tt.splat %xnumel : i32 -> tensor<128x1xi32> loc(#loc42)
+    %xmask_7 = arith.cmpi slt, %xindex_6, %xmask : tensor<128x1xi32> loc(#loc42)
+    %r0_index = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc43)
+    %r0_index_8 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<32xi32> -> tensor<1x32xi32> loc(#loc44)
+    %r0_offset = arith.constant 0 : i32 loc(#loc45)
+    %r0_mask = arith.constant true loc(#loc46)
+    %r0_mask_9 = arith.constant dense<true> : tensor<128x32xi1> loc(#loc46)
+    %tmp0 = arith.extsi %r0_index_8 : tensor<1x32xi32> to tensor<1x32xi64> loc(#loc47)
+    %tmp0_10 = tt.splat %ks0 : i64 -> tensor<1x32xi64> loc(#loc47)
+    %tmp0_11 = arith.muli %tmp0_10, %tmp0 : tensor<1x32xi64> loc(#loc47)
+    %tmp0_12 = arith.extsi %xindex_6 : tensor<128x1xi32> to tensor<128x1xi64> loc(#loc48)
+    %tmp0_13 = tt.broadcast %tmp0_12 : tensor<128x1xi64> -> tensor<128x32xi64> loc(#loc48)
+    %tmp0_14 = tt.broadcast %tmp0_11 : tensor<1x32xi64> -> tensor<128x32xi64> loc(#loc48)
+    %tmp0_15 = arith.addi %tmp0_13, %tmp0_14 : tensor<128x32xi64> loc(#loc48)
+    %tmp0_16 = tt.splat %in_ptr0 : !tt.ptr<f32> -> tensor<128x32x!tt.ptr<f32>> loc(#loc49)
+    %tmp0_17 = tt.addptr %tmp0_16, %tmp0_15 : tensor<128x32x!tt.ptr<f32>>, tensor<128x32xi64> loc(#loc49)
+    %tmp0_18 = arith.constant 0.000000e+00 : f32 loc(#loc50)
+    %tmp0_19 = tt.broadcast %xmask_7 : tensor<128x1xi1> -> tensor<128x32xi1> loc(#loc50)
+    %tmp0_20 = arith.constant dense<0.000000e+00> : tensor<128x32xf32> loc(#loc50)
+    %tmp0_21 = tt.load %tmp0_17, %tmp0_19, %tmp0_20 : tensor<128x32x!tt.ptr<f32>> loc(#loc50)
+    %tmp3 = arith.constant 0 : i32 loc(#loc51)
+    %tmp3_22 = arith.constant 0.000000e+00 : f32 loc(#loc51)
+    %tmp3_23 = arith.constant dense<0.000000e+00> : tensor<128x32xf32> loc(#loc51)
+    %tmp3_24 = tt.broadcast %xmask_7 : tensor<128x1xi1> -> tensor<128x32xi1> loc(#loc51)
+    %tmp3_25 = arith.select %tmp3_24, %tmp0_21, %tmp3_23 : tensor<128x32xi1>, tensor<128x32xf32> loc(#loc51)
+    %tmp4 = tt.call @"triton.language.standard.sum__fp32S128_32S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%tmp3_25) : (tensor<128x32xf32>) -> tensor<128xf32> loc(#loc52)
+    %tmp4_26 = tt.expand_dims %tmp4 {axis = 1 : i32} : tensor<128xf32> -> tensor<128x1xf32> loc(#loc53)
+    %0 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<128x1x!tt.ptr<bf16>> loc(#loc19)
+    %1 = tt.addptr %0, %xindex_6 : tensor<128x1x!tt.ptr<bf16>>, tensor<128x1xi32> loc(#loc19)
+    %2 = arith.truncf %tmp4_26 : tensor<128x1xf32> to tensor<128x1xbf16> loc(#loc20)
+    tt.store %1, %2, %xmask_7 : tensor<128x1x!tt.ptr<bf16>> loc(#loc20)
+    tt.return loc(#loc21)
+  } loc(#loc)
+  tt.func private @"triton.language.standard.sum__fp32S128_32S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%input: tensor<128x32xf32> loc("input"(#loc22))) -> tensor<128xf32> attributes {noinline = false} {
+    %0 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
+    ^bb0(%arg1: f32 loc(unknown), %arg2: f32 loc(unknown)):
+      %2 = tt.call @triton.language.standard._sum_combine__fp32_fp32__(%arg1, %arg2) : (f32, f32) -> f32 loc(#loc23)
+      tt.reduce.return %2 : f32 loc(#loc23)
+    }) : (tensor<128x32xf32>) -> tensor<128xf32> loc(#loc23)
+    tt.return %0 : tensor<128xf32> loc(#loc25)
+  ^bb1:  // no predecessors
+    %1 = ub.poison : tensor<128xf32> loc(#loc26)
+    tt.return %1 : tensor<128xf32> loc(#loc26)
+  } loc(#loc22)
+  tt.func private @triton.language.standard._sum_combine__fp32_fp32__(%a: f32 loc("a"(#loc27)), %b: f32 loc("b"(#loc27))) -> f32 attributes {noinline = false} {
+    %0 = arith.addf %a, %b : f32 loc(#loc28)
+    tt.return %0 : f32 loc(#loc29)
+  ^bb1:  // no predecessors
+    %1 = ub.poison : f32 loc(#loc30)
+    tt.return %1 : f32 loc(#loc30)
+  } loc(#loc27)
+} loc(#loc)
+#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":19:15)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":23:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":23:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:36)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:44)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:23)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":25:21)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":26:28)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":26:38)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":27:16)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":28:48)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:39)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:35)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:30)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:46)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":35:33)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":36:24)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":36:27)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:25)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:36)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:4)
+#loc23 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc25 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:11)
+#loc26 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:4)
+#loc28 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc29 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:11)
+#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:4)
+#loc36 = loc("r0_numel"(#loc1))
+#loc37 = loc("xoffset"(#loc2))
+#loc38 = loc("xoffset"(#loc3))
+#loc39 = loc("xindex"(#loc4))
+#loc40 = loc("xindex"(#loc5))
+#loc41 = loc("xindex"(#loc6))
+#loc42 = loc("xmask"(#loc7))
+#loc43 = loc("r0_index"(#loc8))
+#loc44 = loc("r0_index"(#loc9))
+#loc45 = loc("r0_offset"(#loc10))
+#loc46 = loc("r0_mask"(#loc11))
+#loc47 = loc("tmp0"(#loc12))
+#loc48 = loc("tmp0"(#loc13))
+#loc49 = loc("tmp0"(#loc14))
+#loc50 = loc("tmp0"(#loc15))
+#loc51 = loc("tmp3"(#loc16))
+#loc52 = loc("tmp4"(#loc17))
+#loc53 = loc("tmp4"(#loc18))

SpecForge-ext/cache/compiled_kernels/triton/6/7AX37JF3GC3476WMWTYVB3NGNRIS5J7SJEWETBMPGJB4H67EAVCA/triton_per_fused__to_copy_mul_sum_1.ttgir

@@ -0,0 +1,81 @@
+#blocked = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [4, 2], order = [0, 1]}>
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":18:0)
+#loc1 = loc(unknown)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":36:24)
+#loc20 = loc("in_ptr0"(#loc))
+#loc21 = loc("out_ptr0"(#loc))
+#loc22 = loc("ks0"(#loc))
+#loc23 = loc("xnumel"(#loc))
+#loc24 = loc("r0_numel"(#loc))
+#loc36 = loc("tmp4"(#loc14))
+#loc39 = loc(callsite(#loc1 at #loc36))
+module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 8 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
+  tt.func public @triton_per_fused__to_copy_mul_sum_1(%in_ptr0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %xnumel: i32 loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<0.000000e+00> : tensor<128x32xf32, #blocked> loc(#loc1)
+    %c128_i32 = arith.constant 128 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc25)
+    %xoffset_0 = arith.muli %xoffset, %c128_i32 : i32 loc(#loc26)
+    %xindex = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc27)
+    %xindex_1 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<128xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<128x1xi32, #blocked> loc(#loc27)
+    %xindex_2 = tt.splat %xoffset_0 : i32 -> tensor<128x1xi32, #blocked> loc(#loc28)
+    %xindex_3 = arith.addi %xindex_2, %xindex_1 : tensor<128x1xi32, #blocked> loc(#loc28)
+    %xmask = tt.splat %xnumel : i32 -> tensor<128x1xi32, #blocked> loc(#loc29)
+    %xmask_4 = arith.cmpi slt, %xindex_3, %xmask : tensor<128x1xi32, #blocked> loc(#loc29)
+    %r0_index = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc30)
+    %r0_index_5 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x32xi32, #blocked> loc(#loc30)
+    %tmp0 = arith.extsi %r0_index_5 : tensor<1x32xi32, #blocked> to tensor<1x32xi64, #blocked> loc(#loc31)
+    %tmp0_6 = tt.splat %ks0 : i64 -> tensor<1x32xi64, #blocked> loc(#loc31)
+    %tmp0_7 = arith.muli %tmp0_6, %tmp0 : tensor<1x32xi64, #blocked> loc(#loc31)
+    %tmp0_8 = arith.extsi %xindex_3 : tensor<128x1xi32, #blocked> to tensor<128x1xi64, #blocked> loc(#loc32)
+    %tmp0_9 = tt.broadcast %tmp0_8 : tensor<128x1xi64, #blocked> -> tensor<128x32xi64, #blocked> loc(#loc32)
+    %tmp0_10 = tt.broadcast %tmp0_7 : tensor<1x32xi64, #blocked> -> tensor<128x32xi64, #blocked> loc(#loc32)
+    %tmp0_11 = arith.addi %tmp0_9, %tmp0_10 : tensor<128x32xi64, #blocked> loc(#loc32)
+    %tmp0_12 = tt.splat %in_ptr0 : !tt.ptr<f32> -> tensor<128x32x!tt.ptr<f32>, #blocked> loc(#loc33)
+    %tmp0_13 = tt.addptr %tmp0_12, %tmp0_11 : tensor<128x32x!tt.ptr<f32>, #blocked>, tensor<128x32xi64, #blocked> loc(#loc33)
+    %tmp0_14 = tt.broadcast %xmask_4 : tensor<128x1xi1, #blocked> -> tensor<128x32xi1, #blocked> loc(#loc34)
+    %tmp0_15 = tt.load %tmp0_13, %tmp0_14, %cst : tensor<128x32x!tt.ptr<f32>, #blocked> loc(#loc34)
+    %tmp3 = arith.select %tmp0_14, %tmp0_15, %cst : tensor<128x32xi1, #blocked>, tensor<128x32xf32, #blocked> loc(#loc35)
+    %tmp4 = "tt.reduce"(%tmp3) <{axis = 1 : i32}> ({
+    ^bb0(%tmp4_17: f32 loc(callsite(#loc1 at #loc36)), %tmp4_18: f32 loc(callsite(#loc1 at #loc36))):
+      %tmp4_19 = arith.addf %tmp4_17, %tmp4_18 : f32 loc(#loc40)
+      tt.reduce.return %tmp4_19 : f32 loc(#loc38)
+    }) : (tensor<128x32xf32, #blocked>) -> tensor<128xf32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc38)
+    %tmp4_16 = tt.expand_dims %tmp4 {axis = 1 : i32} : tensor<128xf32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<128x1xf32, #blocked> loc(#loc37)
+    %0 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<128x1x!tt.ptr<bf16>, #blocked> loc(#loc17)
+    %1 = tt.addptr %0, %xindex_3 : tensor<128x1x!tt.ptr<bf16>, #blocked>, tensor<128x1xi32, #blocked> loc(#loc17)
+    %2 = arith.truncf %tmp4_16 : tensor<128x1xf32, #blocked> to tensor<128x1xbf16, #blocked> loc(#loc18)
+    tt.store %1, %2, %xmask_4 : tensor<128x1x!tt.ptr<bf16>, #blocked> loc(#loc18)
+    tt.return loc(#loc19)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":23:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":23:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:44)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":25:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":26:38)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:39)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:35)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:30)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:46)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":35:33)
+#loc13 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc15 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":36:27)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:25)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:36)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:4)
+#loc25 = loc("xoffset"(#loc2))
+#loc26 = loc("xoffset"(#loc3))
+#loc27 = loc("xindex"(#loc4))
+#loc28 = loc("xindex"(#loc5))
+#loc29 = loc("xmask"(#loc6))
+#loc30 = loc("r0_index"(#loc7))
+#loc31 = loc("tmp0"(#loc8))
+#loc32 = loc("tmp0"(#loc9))
+#loc33 = loc("tmp0"(#loc10))
+#loc34 = loc("tmp0"(#loc11))
+#loc35 = loc("tmp3"(#loc12))
+#loc37 = loc("tmp4"(#loc16))
+#loc38 = loc(callsite(#loc13 at #loc36))
+#loc40 = loc(callsite(#loc15 at #loc38))

SpecForge-ext/cache/compiled_kernels/triton/6/7AX37JF3GC3476WMWTYVB3NGNRIS5J7SJEWETBMPGJB4H67EAVCA/triton_per_fused__to_copy_mul_sum_1.ttir

@@ -0,0 +1,84 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":18:0)
+#loc1 = loc(unknown)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":36:24)
+#loc22 = loc("in_ptr0"(#loc))
+#loc23 = loc("out_ptr0"(#loc))
+#loc24 = loc("ks0"(#loc))
+#loc25 = loc("xnumel"(#loc))
+#loc26 = loc("r0_numel"(#loc))
+#loc40 = loc("tmp4"(#loc16))
+#loc43 = loc(callsite(#loc1 at #loc40))
+module {
+  tt.func public @triton_per_fused__to_copy_mul_sum_1(%in_ptr0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %xnumel: i32 loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<0.000000e+00> : tensor<128x32xf32> loc(#loc1)
+    %c128_i32 = arith.constant 128 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc27)
+    %xoffset_0 = arith.muli %xoffset, %c128_i32 : i32 loc(#loc28)
+    %xindex = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc29)
+    %xindex_1 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<128xi32> -> tensor<128x1xi32> loc(#loc30)
+    %xindex_2 = tt.splat %xoffset_0 : i32 -> tensor<128x1xi32> loc(#loc31)
+    %xindex_3 = arith.addi %xindex_2, %xindex_1 : tensor<128x1xi32> loc(#loc31)
+    %xmask = tt.splat %xnumel : i32 -> tensor<128x1xi32> loc(#loc32)
+    %xmask_4 = arith.cmpi slt, %xindex_3, %xmask : tensor<128x1xi32> loc(#loc32)
+    %r0_index = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc33)
+    %r0_index_5 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<32xi32> -> tensor<1x32xi32> loc(#loc34)
+    %tmp0 = arith.extsi %r0_index_5 : tensor<1x32xi32> to tensor<1x32xi64> loc(#loc35)
+    %tmp0_6 = tt.splat %ks0 : i64 -> tensor<1x32xi64> loc(#loc35)
+    %tmp0_7 = arith.muli %tmp0_6, %tmp0 : tensor<1x32xi64> loc(#loc35)
+    %tmp0_8 = arith.extsi %xindex_3 : tensor<128x1xi32> to tensor<128x1xi64> loc(#loc36)
+    %tmp0_9 = tt.broadcast %tmp0_8 : tensor<128x1xi64> -> tensor<128x32xi64> loc(#loc36)
+    %tmp0_10 = tt.broadcast %tmp0_7 : tensor<1x32xi64> -> tensor<128x32xi64> loc(#loc36)
+    %tmp0_11 = arith.addi %tmp0_9, %tmp0_10 : tensor<128x32xi64> loc(#loc36)
+    %tmp0_12 = tt.splat %in_ptr0 : !tt.ptr<f32> -> tensor<128x32x!tt.ptr<f32>> loc(#loc37)
+    %tmp0_13 = tt.addptr %tmp0_12, %tmp0_11 : tensor<128x32x!tt.ptr<f32>>, tensor<128x32xi64> loc(#loc37)
+    %tmp0_14 = tt.broadcast %xmask_4 : tensor<128x1xi1> -> tensor<128x32xi1> loc(#loc38)
+    %tmp0_15 = tt.load %tmp0_13, %tmp0_14, %cst : tensor<128x32x!tt.ptr<f32>> loc(#loc38)
+    %tmp3 = arith.select %tmp0_14, %tmp0_15, %cst : tensor<128x32xi1>, tensor<128x32xf32> loc(#loc39)
+    %tmp4 = "tt.reduce"(%tmp3) <{axis = 1 : i32}> ({
+    ^bb0(%tmp4_17: f32 loc(callsite(#loc1 at #loc40)), %tmp4_18: f32 loc(callsite(#loc1 at #loc40))):
+      %tmp4_19 = arith.addf %tmp4_17, %tmp4_18 : f32 loc(#loc44)
+      tt.reduce.return %tmp4_19 : f32 loc(#loc42)
+    }) : (tensor<128x32xf32>) -> tensor<128xf32> loc(#loc42)
+    %tmp4_16 = tt.expand_dims %tmp4 {axis = 1 : i32} : tensor<128xf32> -> tensor<128x1xf32> loc(#loc41)
+    %0 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<128x1x!tt.ptr<bf16>> loc(#loc19)
+    %1 = tt.addptr %0, %xindex_3 : tensor<128x1x!tt.ptr<bf16>>, tensor<128x1xi32> loc(#loc19)
+    %2 = arith.truncf %tmp4_16 : tensor<128x1xf32> to tensor<128x1xbf16> loc(#loc20)
+    tt.store %1, %2, %xmask_4 : tensor<128x1x!tt.ptr<bf16>> loc(#loc20)
+    tt.return loc(#loc21)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":23:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":23:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:36)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:44)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":24:23)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":25:21)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":26:28)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":26:38)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:39)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:35)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:30)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":33:46)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":35:33)
+#loc15 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc17 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":36:27)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:25)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:36)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/yf/cyft7sialepriw6eujulaxpi57qlrafkmp4k2kjwzw4noh23ddz6.py":37:4)
+#loc27 = loc("xoffset"(#loc2))
+#loc28 = loc("xoffset"(#loc3))
+#loc29 = loc("xindex"(#loc4))
+#loc30 = loc("xindex"(#loc5))
+#loc31 = loc("xindex"(#loc6))
+#loc32 = loc("xmask"(#loc7))
+#loc33 = loc("r0_index"(#loc8))
+#loc34 = loc("r0_index"(#loc9))
+#loc35 = loc("tmp0"(#loc10))
+#loc36 = loc("tmp0"(#loc11))
+#loc37 = loc("tmp0"(#loc12))
+#loc38 = loc("tmp0"(#loc13))
+#loc39 = loc("tmp3"(#loc14))
+#loc41 = loc("tmp4"(#loc18))
+#loc42 = loc(callsite(#loc15 at #loc40))
+#loc44 = loc(callsite(#loc17 at #loc42))

SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/__grp__triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.source", "triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ttir", "triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ttgir", "triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.llir", "triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ptx", "triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.cubin", "triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.json"}}

SpecForge-ext/cache/compiled_kernels/triton/6/C3FCZCDEMCLSFODWXLEH5MRAQRWLOTRP4SAQURVAE7BPHZSTV2WQ/triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0.ttir

@@ -0,0 +1,233 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":18:0)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":49:33)
+#loc3 = loc(unknown)
+#loc35 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":178:28)
+#loc42 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":181:58)
+#loc59 = loc("in_ptr0"(#loc))
+#loc60 = loc("out_ptr2"(#loc))
+#loc61 = loc("xnumel"(#loc))
+#loc62 = loc("r0_numel"(#loc))
+#loc90 = loc("out_max"(#loc35))
+#loc96 = loc("out_sum"(#loc42))
+#loc123 = loc(callsite(#loc90 at #loc2))
+#loc129 = loc(callsite(#loc96 at #loc2))
+#loc138 = loc(callsite(#loc3 at #loc123))
+#loc141 = loc(callsite(#loc3 at #loc129))
+module {
+  tt.func public @triton_red_fused__softmax__to_copy_exp_prepare_softmax_online_sub_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr2"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %delta = arith.constant dense<0xFF800000> : tensor<1x1xf32> loc(#loc108)
+    %cst = arith.constant dense<1.000000e+00> : tensor<1x2048xf32> loc(#loc3)
+    %cst_0 = arith.constant dense<0.000000e+00> : tensor<1x2048xf32> loc(#loc3)
+    %cst_1 = arith.constant dense<0.000000e+00> : tensor<1x2048xbf16> loc(#loc3)
+    %c2048_i32 = arith.constant 2048 : i32 loc(#loc3)
+    %c32000_i32 = arith.constant 32000 : i32 loc(#loc3)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc3)
+    %cst_2 = arith.constant dense<32000> : tensor<1x2048xi32> loc(#loc3)
+    %cst_3 = arith.constant dense<0xFF800000> : tensor<1x2048xf32> loc(#loc3)
+    %xoffset = tt.get_program_id x : i32 loc(#loc64)
+    %r0_base = tt.make_range {end = 2048 : i32, start = 0 : i32} : tensor<2048xi32> loc(#loc65)
+    %r0_base_4 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<2048xi32> -> tensor<1x2048xi32> loc(#loc66)
+    %_tmp3_sum:2 = scf.for %r0_offset = %c0_i32 to %c32000_i32 step %c2048_i32 iter_args(%_tmp3_max = %cst_3, %_tmp3_sum_12 = %cst_0) -> (tensor<1x2048xf32>, tensor<1x2048xf32>)  : i32 {
+      %r0_index = tt.splat %r0_offset : i32 -> tensor<1x2048xi32> loc(#loc68)
+      %r0_index_13 = arith.addi %r0_index, %r0_base_4 : tensor<1x2048xi32> loc(#loc68)
+      %r0_mask = arith.cmpi slt, %r0_index_13, %cst_2 : tensor<1x2048xi32> loc(#loc69)
+      %tmp0 = arith.muli %xoffset, %c32000_i32 : i32 loc(#loc70)
+      %tmp0_14 = tt.splat %tmp0 : i32 -> tensor<1x2048xi32> loc(#loc110)
+      %tmp0_15 = arith.addi %r0_index_13, %tmp0_14 : tensor<1x2048xi32> loc(#loc71)
+      %tmp0_16 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1x2048x!tt.ptr<bf16>> loc(#loc72)
+      %tmp0_17 = tt.addptr %tmp0_16, %tmp0_15 : tensor<1x2048x!tt.ptr<bf16>>, tensor<1x2048xi32> loc(#loc72)
+      %tmp0_18 = tt.load %tmp0_17, %r0_mask, %cst_1 evictionPolicy = evict_last : tensor<1x2048x!tt.ptr<bf16>> loc(#loc73)
+      %tmp0_19 = arith.extf %tmp0_18 : tensor<1x2048xbf16> to tensor<1x2048xf32> loc(#loc74)
+      %mask = arith.cmpf ogt, %_tmp3_max, %tmp0_19 : tensor<1x2048xf32> loc(#loc131)
+      %mask_20 = arith.cmpf une, %_tmp3_max, %_tmp3_max : tensor<1x2048xf32> loc(#loc132)
+      %mask_21 = arith.ori %mask, %mask_20 : tensor<1x2048xi1> loc(#loc133)
+      %out_max_22 = arith.select %mask_21, %_tmp3_max, %tmp0_19 : tensor<1x2048xi1>, tensor<1x2048xf32> loc(#loc134)
+      %lhs_scale = arith.cmpf oeq, %out_max_22, %cst_3 : tensor<1x2048xf32> loc(#loc114)
+      %lhs_scale_23 = arith.subf %_tmp3_max, %out_max_22 : tensor<1x2048xf32> loc(#loc115)
+      %lhs_scale_24 = tt.extern_elementwise %lhs_scale_23 {libname = "", libpath = "", pure = true, symbol = "__nv_expf"} : (tensor<1x2048xf32>) -> tensor<1x2048xf32> loc(#loc135)
+      %lhs_scale_25 = arith.select %lhs_scale, %cst, %lhs_scale_24 : tensor<1x2048xi1>, tensor<1x2048xf32> loc(#loc117)
+      %rhs_scale = arith.subf %tmp0_19, %out_max_22 : tensor<1x2048xf32> loc(#loc118)
+      %rhs_scale_26 = tt.extern_elementwise %rhs_scale {libname = "", libpath = "", pure = true, symbol = "__nv_expf"} : (tensor<1x2048xf32>) -> tensor<1x2048xf32> loc(#loc136)
+      %rhs_scale_27 = arith.select %lhs_scale, %cst, %rhs_scale_26 : tensor<1x2048xi1>, tensor<1x2048xf32> loc(#loc120)
+      %out_sum_28 = arith.mulf %_tmp3_sum_12, %lhs_scale_25 : tensor<1x2048xf32> loc(#loc121)
+      %out_sum_29 = arith.addf %out_sum_28, %rhs_scale_27 : tensor<1x2048xf32> loc(#loc122)
+      %_tmp3_max_30 = arith.select %r0_mask, %out_max_22, %_tmp3_max : tensor<1x2048xi1>, tensor<1x2048xf32> loc(#loc88)
+      %_tmp3_sum_31 = arith.select %r0_mask, %out_sum_29, %_tmp3_sum_12 : tensor<1x2048xi1>, tensor<1x2048xf32> loc(#loc89)
+      scf.yield %_tmp3_max_30, %_tmp3_sum_31 : tensor<1x2048xf32>, tensor<1x2048xf32> loc(#loc33)
+    } loc(#loc109)
+    %out_max = "tt.reduce"(%_tmp3_sum#0) <{axis = 1 : i32}> ({
+    ^bb0(%out_max_12: f32 loc(callsite(#loc3 at #loc123)), %out_max_13: f32 loc(callsite(#loc3 at #loc123))):
+      %mask = arith.cmpf ogt, %out_max_12, %out_max_13 : f32 loc(#loc142)
+      %mask_14 = arith.cmpf une, %out_max_12, %out_max_12 : f32 loc(#loc143)
+      %mask_15 = arith.ori %mask, %mask_14 : i1 loc(#loc144)
+      %out_max_16 = arith.select %mask_15, %out_max_12, %out_max_13 : f32 loc(#loc145)
+      tt.reduce.return %out_max_16 : f32 loc(#loc137)
+    }) : (tensor<1x2048xf32>) -> tensor<1xf32> loc(#loc137)
+    %out_max_keepdim = tt.expand_dims %out_max {axis = 1 : i32} : tensor<1xf32> -> tensor<1x1xf32> loc(#loc124)
+    %delta_5 = arith.cmpf oeq, %out_max_keepdim, %delta : tensor<1x1xf32> loc(#loc108)
+    %delta_6 = tt.broadcast %out_max_keepdim : tensor<1x1xf32> -> tensor<1x2048xf32> loc(#loc125)
+    %delta_7 = arith.subf %_tmp3_sum#0, %delta_6 : tensor<1x2048xf32> loc(#loc125)
+    %delta_8 = tt.broadcast %delta_5 : tensor<1x1xi1> -> tensor<1x2048xi1> loc(#loc126)
+    %delta_9 = arith.select %delta_8, %cst_0, %delta_7 : tensor<1x2048xi1>, tensor<1x2048xf32> loc(#loc126)
+    %out_sum = tt.extern_elementwise %delta_9 {libname = "", libpath = "", pure = true, symbol = "__nv_expf"} : (tensor<1x2048xf32>) -> tensor<1x2048xf32> loc(#loc139)
+    %out_sum_10 = arith.mulf %_tmp3_sum#1, %out_sum : tensor<1x2048xf32> loc(#loc128)
+    %out_sum_11 = "tt.reduce"(%out_sum_10) <{axis = 1 : i32}> ({
+    ^bb0(%out_sum_12: f32 loc(callsite(#loc3 at #loc129)), %out_sum_13: f32 loc(callsite(#loc3 at #loc129))):
+      %out_sum_14 = arith.addf %out_sum_12, %out_sum_13 : f32 loc(#loc146)
+      tt.reduce.return %out_sum_14 : f32 loc(#loc140)
+    }) : (tensor<1x2048xf32>) -> tensor<1xf32> loc(#loc140)
+    %tmp4 = tt.expand_dims %out_sum_11 {axis = 1 : i32} : tensor<1xf32> -> tensor<1x1xf32> loc(#loc97)
+    scf.for %r0_offset = %c0_i32 to %c32000_i32 step %c2048_i32  : i32 {
+      %r0_index = tt.splat %r0_offset : i32 -> tensor<1x2048xi32> loc(#loc98)
+      %r0_index_12 = arith.addi %r0_index, %r0_base_4 : tensor<1x2048xi32> loc(#loc98)
+      %r0_mask = arith.cmpi slt, %r0_index_12, %cst_2 : tensor<1x2048xi32> loc(#loc99)
+      %tmp5 = arith.muli %xoffset, %c32000_i32 : i32 loc(#loc100)
+      %tmp5_13 = tt.splat %tmp5 : i32 -> tensor<1x2048xi32> loc(#loc130)
+      %tmp5_14 = arith.addi %r0_index_12, %tmp5_13 : tensor<1x2048xi32> loc(#loc101)
+      %tmp5_15 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1x2048x!tt.ptr<bf16>> loc(#loc102)
+      %tmp5_16 = tt.addptr %tmp5_15, %tmp5_14 : tensor<1x2048x!tt.ptr<bf16>>, tensor<1x2048xi32> loc(#loc102)
+      %tmp5_17 = tt.load %tmp5_16, %r0_mask, %cst_1 evictionPolicy = evict_first : tensor<1x2048x!tt.ptr<bf16>> loc(#loc103)
+      %tmp5_18 = arith.extf %tmp5_17 : tensor<1x2048xbf16> to tensor<1x2048xf32> loc(#loc104)
+      %tmp7 = arith.subf %tmp5_18, %delta_6 : tensor<1x2048xf32> loc(#loc105)
+      %tmp8 = tt.extern_elementwise %tmp7 {libname = "", libpath = "", pure = true, symbol = "__nv_expf"} : (tensor<1x2048xf32>) -> tensor<1x2048xf32> loc(#loc106)
+      %tmp9 = tt.broadcast %tmp4 : tensor<1x1xf32> -> tensor<1x2048xf32> loc(#loc107)
+      %tmp9_19 = arith.divf %tmp8, %tmp9 : tensor<1x2048xf32> loc(#loc107)
+      %0 = tt.splat %out_ptr2 : !tt.ptr<f32> -> tensor<1x2048x!tt.ptr<f32>> loc(#loc56)
+      %1 = tt.addptr %0, %tmp5_14 : tensor<1x2048x!tt.ptr<f32>>, tensor<1x2048xi32> loc(#loc56)
+      tt.store %1, %tmp9_19, %r0_mask : tensor<1x2048x!tt.ptr<f32>> loc(#loc57)
+    } loc(#loc45)
+    tt.return loc(#loc58)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":180:40)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":23:28)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":26:27)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":26:37)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":31:40)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":32:31)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":33:29)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":37:47)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":37:41)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":37:34)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":37:52)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":37:105)
+#loc15 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":110:15)
+#loc16 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":193:31)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":42:40)
+#loc18 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:21)
+#loc19 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:16)
+#loc20 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":113:29)
+#loc21 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":196:19)
+#loc22 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":196:53)
+#loc23 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":173:29)
+#loc24 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":196:62)
+#loc25 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":196:39)
+#loc26 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":199:53)
+#loc27 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":199:62)
+#loc28 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":199:39)
+#loc29 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":205:24)
+#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":205:36)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":45:54)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":46:54)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":46:8)
+#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":123:29)
+#loc36 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":179:46)
+#loc37 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":180:68)
+#loc38 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":180:58)
+#loc39 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":181:42)
+#loc40 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":181:31)
+#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc43 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":51:16)
+#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":52:40)
+#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":53:31)
+#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":54:29)
+#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":58:47)
+#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":58:41)
+#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":58:34)
+#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":58:52)
+#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":58:106)
+#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":60:22)
+#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":61:29)
+#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":62:23)
+#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":63:29)
+#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":63:53)
+#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/cz/cczmmrlu3u67jn3gitfwyu3jif3b3pwoxj37xp4ddackcufwiqri.py":52:4)
+#loc63 = loc("delta"(#loc1))
+#loc64 = loc("xoffset"(#loc4))
+#loc65 = loc("r0_base"(#loc5))
+#loc66 = loc("r0_base"(#loc6))
+#loc67 = loc("_tmp3_max"(#loc7))
+#loc68 = loc("r0_index"(#loc8))
+#loc69 = loc("r0_mask"(#loc9))
+#loc70 = loc("tmp0"(#loc10))
+#loc71 = loc("tmp0"(#loc11))
+#loc72 = loc("tmp0"(#loc12))
+#loc73 = loc("tmp0"(#loc13))
+#loc74 = loc("tmp0"(#loc14))
+#loc75 = loc("mask"(#loc15))
+#loc76 = loc("out_max"(#loc16))
+#loc77 = loc("mask"(#loc18))
+#loc78 = loc("mask"(#loc19))
+#loc79 = loc("lhs_scale"(#loc21))
+#loc80 = loc("lhs_scale"(#loc22))
+#loc81 = loc("lhs_scale"(#loc24))
+#loc82 = loc("lhs_scale"(#loc25))
+#loc83 = loc("rhs_scale"(#loc26))
+#loc84 = loc("rhs_scale"(#loc27))
+#loc85 = loc("rhs_scale"(#loc28))
+#loc86 = loc("out_sum"(#loc29))
+#loc87 = loc("out_sum"(#loc30))
+#loc88 = loc("_tmp3_max"(#loc31))
+#loc89 = loc("_tmp3_sum"(#loc32))
+#loc91 = loc("out_max_keepdim"(#loc36))
+#loc92 = loc("delta"(#loc37))
+#loc93 = loc("delta"(#loc38))
+#loc94 = loc("out_sum"(#loc39))
+#loc95 = loc("out_sum"(#loc40))
+#loc97 = loc("tmp4"(#loc44))
+#loc98 = loc("r0_index"(#loc46))
+#loc99 = loc("r0_mask"(#loc47))
+#loc100 = loc("tmp5"(#loc48))
+#loc101 = loc("tmp5"(#loc49))
+#loc102 = loc("tmp5"(#loc50))
+#loc103 = loc("tmp5"(#loc51))
+#loc104 = loc("tmp5"(#loc52))
+#loc105 = loc("tmp7"(#loc53))
+#loc106 = loc("tmp8"(#loc54))
+#loc107 = loc("tmp9"(#loc55))
+#loc108 = loc(callsite(#loc63 at #loc2))
+#loc109 = loc("_tmp3_sum"(#loc67))
+#loc110 = loc(fused[#loc71, #loc70])
+#loc111 = loc("mask"(#loc75))
+#loc112 = loc(callsite(#loc76 at #loc17))
+#loc113 = loc("mask"(#loc78))
+#loc114 = loc(callsite(#loc79 at #loc17))
+#loc115 = loc(callsite(#loc80 at #loc17))
+#loc116 = loc(callsite(#loc81 at #loc17))
+#loc117 = loc(callsite(#loc82 at #loc17))
+#loc118 = loc(callsite(#loc83 at #loc17))
+#loc119 = loc(callsite(#loc84 at #loc17))
+#loc120 = loc(callsite(#loc85 at #loc17))
+#loc121 = loc(callsite(#loc86 at #loc17))
+#loc122 = loc(callsite(#loc87 at #loc17))
+#loc124 = loc(callsite(#loc91 at #loc2))
+#loc125 = loc(callsite(#loc92 at #loc2))
+#loc126 = loc(callsite(#loc93 at #loc2))
+#loc127 = loc(callsite(#loc94 at #loc2))
+#loc128 = loc(callsite(#loc95 at #loc2))
+#loc130 = loc(fused[#loc101, #loc100])
+#loc131 = loc(callsite(#loc111 at #loc112))
+#loc132 = loc(callsite(#loc77 at #loc112))
+#loc133 = loc(callsite(#loc113 at #loc112))
+#loc134 = loc(callsite(#loc20 at #loc112))
+#loc135 = loc(callsite(#loc23 at #loc116))
+#loc136 = loc(callsite(#loc23 at #loc119))
+#loc137 = loc(callsite(#loc34 at #loc123))
+#loc139 = loc(callsite(#loc23 at #loc127))
+#loc140 = loc(callsite(#loc41 at #loc129))
+#loc142 = loc(callsite(#loc111 at #loc137))
+#loc143 = loc(callsite(#loc77 at #loc137))
+#loc144 = loc(callsite(#loc113 at #loc137))
+#loc145 = loc(callsite(#loc20 at #loc137))
+#loc146 = loc(callsite(#loc43 at #loc140))

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/__grp__triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.source", "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ttir", "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ttgir", "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.llir", "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ptx", "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.cubin", "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.json"}}

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.json

@@ -0,0 +1 @@
+{"hash": "2d88f34765cb5a3f781885efc3427f98913e07cbedc15d0af342925a3b9998ab", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 4, "num_ctas": 1, "num_stages": 1, "warp_size": 32, "maxnreg": null, "cluster_dims": [1, 1, 1], "ptx_version": null, "ptx_options": null, "ir_override": null, "enable_fp_fusion": true, "launch_cooperative_grid": false, "launch_pdl": false, "supported_fp8_dtypes": ["fp8e4b15", "fp8e4nv", "fp8e5"], "deprecated_fp8_dot_operand_dtypes": ["fp8e4b15"], "default_dot_input_precision": "tf32", "allowed_dot_input_precisions": ["tf32", "tf32x3", "ieee"], "max_num_imprecise_acc_default": 1073741824, "extern_libs": [["libdevice", "/workspace/specforge/lib/python3.11/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]], "debug": true, "backend_name": "cuda", "sanitize_overflow": false, "arch": "sm90", "instrumentation_mode": "", "triton_version": "3.5.1", "tensordesc_meta": [], "shared": 0, "tmem_size": 0, "global_scratch_size": 0, "global_scratch_align": 1, "profile_scratch_size": 0, "profile_scratch_align": 1, "name": "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0"}

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.llir

@@ -0,0 +1,808 @@
+; ModuleID = 'LLVMDialectModule'
+source_filename = "LLVMDialectModule"
+target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+
+@assertFunc_2 = internal constant [8 x i8] c"unknown\00"
+@assertFile_2 = internal constant [114 x i8] c"/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py\00"
+@assertMessage_2 = internal constant [38 x i8] c"index out of bounds: 0 <= tmp36 < ks3\00"
+@assertFunc_1 = internal constant [8 x i8] c"unknown\00"
+@assertFile_1 = internal constant [114 x i8] c"/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py\00"
+@assertMessage_1 = internal constant [65 x i8] c"index out of bounds: 0 <= tl.broadcast_to(tmp23, [XBLOCK]) < ks2\00"
+@assertFunc_0 = internal constant [8 x i8] c"unknown\00"
+@assertFile_0 = internal constant [114 x i8] c"/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py\00"
+@assertMessage_0 = internal constant [64 x i8] c"index out of bounds: 0 <= tl.broadcast_to(tmp8, [XBLOCK]) < ks2\00"
+
+; Function Attrs: noreturn
+declare !dbg !5 void @__assertfail(ptr, ptr, i32, ptr, i64) local_unnamed_addr #0
+
+define ptx_kernel void @triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, ptr addrspace(1) %4, i64 %5, i64 %6, i64 %7, i64 %8, i32 %9, ptr addrspace(1) readnone captures(none) %10, ptr addrspace(1) readnone captures(none) %11) local_unnamed_addr #1 !dbg !9 {
+  %13 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !10
+  %14 = shl i32 %13, 10, !dbg !11
+  %15 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !12
+  %16 = shl nuw nsw i32 %15, 3, !dbg !12
+  %17 = and i32 %16, 1016, !dbg !12
+  %18 = insertelement <8 x i32> poison, i32 %9, i64 0, !dbg !13
+  %19 = shufflevector <8 x i32> %18, <8 x i32> poison, <8 x i32> zeroinitializer, !dbg !13
+  %20 = or disjoint i32 %17, %14, !dbg !14
+  %21 = or disjoint i32 %20, 1, !dbg !14
+  %22 = or disjoint i32 %20, 2, !dbg !14
+  %23 = or disjoint i32 %20, 3, !dbg !14
+  %24 = or disjoint i32 %20, 4, !dbg !14
+  %25 = or disjoint i32 %20, 5, !dbg !14
+  %26 = or disjoint i32 %20, 6, !dbg !14
+  %27 = or disjoint i32 %20, 7, !dbg !14
+  %28 = insertelement <8 x i32> poison, i32 %20, i64 0, !dbg !13
+  %29 = insertelement <8 x i32> %28, i32 %21, i64 1, !dbg !13
+  %30 = insertelement <8 x i32> %29, i32 %22, i64 2, !dbg !13
+  %31 = insertelement <8 x i32> %30, i32 %23, i64 3, !dbg !13
+  %32 = insertelement <8 x i32> %31, i32 %24, i64 4, !dbg !13
+  %33 = insertelement <8 x i32> %32, i32 %25, i64 5, !dbg !13
+  %34 = insertelement <8 x i32> %33, i32 %26, i64 6, !dbg !13
+  %35 = insertelement <8 x i32> %34, i32 %27, i64 7, !dbg !13
+  %36 = icmp slt <8 x i32> %35, %19, !dbg !13
+  %37 = sext <8 x i32> %35 to <8 x i64>, !dbg !15
+  %38 = insertelement <8 x i64> poison, i64 %5, i64 0, !dbg !15
+  %39 = shufflevector <8 x i64> %38, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !15
+  %40 = srem <8 x i64> %37, %39, !dbg !15
+  %41 = extractelement <8 x i64> %37, i64 0, !dbg !16
+  %42 = sdiv i64 %41, %5, !dbg !17
+  %43 = extractelement <8 x i64> %37, i64 1, !dbg !16
+  %44 = sdiv i64 %43, %5, !dbg !17
+  %45 = extractelement <8 x i64> %37, i64 2, !dbg !16
+  %46 = sdiv i64 %45, %5, !dbg !17
+  %47 = extractelement <8 x i64> %37, i64 3, !dbg !16
+  %48 = sdiv i64 %47, %5, !dbg !17
+  %49 = extractelement <8 x i64> %37, i64 4, !dbg !16
+  %50 = sdiv i64 %49, %5, !dbg !17
+  %51 = extractelement <8 x i64> %37, i64 5, !dbg !16
+  %52 = sdiv i64 %51, %5, !dbg !17
+  %53 = extractelement <8 x i64> %37, i64 6, !dbg !16
+  %54 = sdiv i64 %53, %5, !dbg !17
+  %55 = extractelement <8 x i64> %37, i64 7, !dbg !16
+  %56 = sdiv i64 %55, %5, !dbg !17
+  %57 = srem i64 %42, %6, !dbg !18
+  %58 = srem i64 %44, %6, !dbg !18
+  %59 = srem i64 %46, %6, !dbg !18
+  %60 = srem i64 %48, %6, !dbg !18
+  %61 = srem i64 %50, %6, !dbg !18
+  %62 = srem i64 %52, %6, !dbg !18
+  %63 = srem i64 %54, %6, !dbg !18
+  %64 = srem i64 %56, %6, !dbg !18
+  %65 = getelementptr bfloat, ptr addrspace(1) %0, i64 %41, !dbg !16
+  %66 = getelementptr bfloat, ptr addrspace(1) %0, i64 %43, !dbg !16
+  %67 = getelementptr bfloat, ptr addrspace(1) %0, i64 %45, !dbg !16
+  %68 = getelementptr bfloat, ptr addrspace(1) %0, i64 %47, !dbg !16
+  %69 = getelementptr bfloat, ptr addrspace(1) %0, i64 %49, !dbg !16
+  %70 = getelementptr bfloat, ptr addrspace(1) %0, i64 %51, !dbg !16
+  %71 = getelementptr bfloat, ptr addrspace(1) %0, i64 %53, !dbg !16
+  %72 = getelementptr bfloat, ptr addrspace(1) %0, i64 %55, !dbg !16
+  %73 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %74 = extractelement <8 x i1> %36, i64 0, !dbg !19
+  %75 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %65, i64 %73, i1 %74) #4, !dbg !19
+  %76 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %77 = extractelement <8 x i1> %36, i64 1, !dbg !19
+  %78 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %66, i64 %76, i1 %77) #4, !dbg !19
+  %79 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %80 = extractelement <8 x i1> %36, i64 2, !dbg !19
+  %81 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %67, i64 %79, i1 %80) #4, !dbg !19
+  %82 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %83 = extractelement <8 x i1> %36, i64 3, !dbg !19
+  %84 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %68, i64 %82, i1 %83) #4, !dbg !19
+  %85 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %86 = extractelement <8 x i1> %36, i64 4, !dbg !19
+  %87 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %69, i64 %85, i1 %86) #4, !dbg !19
+  %88 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %89 = extractelement <8 x i1> %36, i64 5, !dbg !19
+  %90 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %70, i64 %88, i1 %89) #4, !dbg !19
+  %91 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %92 = extractelement <8 x i1> %36, i64 6, !dbg !19
+  %93 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %71, i64 %91, i1 %92) #4, !dbg !19
+  %94 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !19
+  %95 = extractelement <8 x i1> %36, i64 7, !dbg !19
+  %96 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %72, i64 %94, i1 %95) #4, !dbg !19
+  %97 = getelementptr i64, ptr addrspace(1) %1, i64 %57, !dbg !20
+  %98 = getelementptr i64, ptr addrspace(1) %1, i64 %58, !dbg !20
+  %99 = getelementptr i64, ptr addrspace(1) %1, i64 %59, !dbg !20
+  %100 = getelementptr i64, ptr addrspace(1) %1, i64 %60, !dbg !20
+  %101 = getelementptr i64, ptr addrspace(1) %1, i64 %61, !dbg !20
+  %102 = getelementptr i64, ptr addrspace(1) %1, i64 %62, !dbg !20
+  %103 = getelementptr i64, ptr addrspace(1) %1, i64 %63, !dbg !20
+  %104 = getelementptr i64, ptr addrspace(1) %1, i64 %64, !dbg !20
+  %105 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %106 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %97, i64 %105, i1 %74) #4, !dbg !21
+  %107 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %108 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %98, i64 %107, i1 %77) #4, !dbg !21
+  %109 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %110 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %99, i64 %109, i1 %80) #4, !dbg !21
+  %111 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %112 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %100, i64 %111, i1 %83) #4, !dbg !21
+  %113 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %114 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %101, i64 %113, i1 %86) #4, !dbg !21
+  %115 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %116 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %102, i64 %115, i1 %89) #4, !dbg !21
+  %117 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %118 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %103, i64 %117, i1 %92) #4, !dbg !21
+  %119 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !21
+  %120 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %104, i64 %119, i1 %95) #4, !dbg !21
+  %121 = sdiv i64 %5, 2, !dbg !22
+  %122 = sub nsw i64 %41, %121, !dbg !23
+  %123 = sub nsw i64 %43, %121, !dbg !23
+  %124 = sub nsw i64 %45, %121, !dbg !23
+  %125 = sub nsw i64 %47, %121, !dbg !23
+  %126 = sub nsw i64 %49, %121, !dbg !23
+  %127 = sub nsw i64 %51, %121, !dbg !23
+  %128 = sub nsw i64 %53, %121, !dbg !23
+  %129 = sub nsw i64 %55, %121, !dbg !23
+  %130 = getelementptr bfloat, ptr addrspace(1) %0, i64 %122, !dbg !24
+  %131 = getelementptr bfloat, ptr addrspace(1) %0, i64 %123, !dbg !24
+  %132 = getelementptr bfloat, ptr addrspace(1) %0, i64 %124, !dbg !24
+  %133 = getelementptr bfloat, ptr addrspace(1) %0, i64 %125, !dbg !24
+  %134 = getelementptr bfloat, ptr addrspace(1) %0, i64 %126, !dbg !24
+  %135 = getelementptr bfloat, ptr addrspace(1) %0, i64 %127, !dbg !24
+  %136 = getelementptr bfloat, ptr addrspace(1) %0, i64 %128, !dbg !24
+  %137 = getelementptr bfloat, ptr addrspace(1) %0, i64 %129, !dbg !24
+  %138 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %139 = insertelement <8 x i64> poison, i64 %121, i64 0, !dbg !26
+  %140 = shufflevector <8 x i64> %139, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !26
+  %141 = icmp sge <8 x i64> %40, %140, !dbg !26
+  %142 = and <8 x i1> %36, %141, !dbg !27
+  %143 = extractelement <8 x i1> %142, i64 0, !dbg !28
+  %144 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %130, i64 %138, i1 %143) #4, !dbg !25
+  %145 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %146 = extractelement <8 x i1> %142, i64 1, !dbg !28
+  %147 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %131, i64 %145, i1 %146) #4, !dbg !25
+  %148 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %149 = extractelement <8 x i1> %142, i64 2, !dbg !28
+  %150 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %132, i64 %148, i1 %149) #4, !dbg !25
+  %151 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %152 = extractelement <8 x i1> %142, i64 3, !dbg !28
+  %153 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %133, i64 %151, i1 %152) #4, !dbg !25
+  %154 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %155 = extractelement <8 x i1> %142, i64 4, !dbg !28
+  %156 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %134, i64 %154, i1 %155) #4, !dbg !25
+  %157 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %158 = extractelement <8 x i1> %142, i64 5, !dbg !28
+  %159 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %135, i64 %157, i1 %158) #4, !dbg !25
+  %160 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %161 = extractelement <8 x i1> %142, i64 6, !dbg !28
+  %162 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %136, i64 %160, i1 %161) #4, !dbg !25
+  %163 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !25
+  %164 = extractelement <8 x i1> %142, i64 7, !dbg !28
+  %165 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %137, i64 %163, i1 %164) #4, !dbg !25
+  %166 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %167 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %97, i64 %166, i1 %143) #4, !dbg !28
+  %168 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %169 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %98, i64 %168, i1 %146) #4, !dbg !28
+  %170 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %171 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %99, i64 %170, i1 %149) #4, !dbg !28
+  %172 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %173 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %100, i64 %172, i1 %152) #4, !dbg !28
+  %174 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %175 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %101, i64 %174, i1 %155) #4, !dbg !28
+  %176 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %177 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %102, i64 %176, i1 %158) #4, !dbg !28
+  %178 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %179 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %103, i64 %178, i1 %161) #4, !dbg !28
+  %180 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !28
+  %181 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %104, i64 %180, i1 %164) #4, !dbg !28
+  %182 = insertelement <8 x i64> poison, i64 %167, i64 0, !dbg !29
+  %183 = insertelement <8 x i64> %182, i64 %169, i64 1, !dbg !29
+  %184 = insertelement <8 x i64> %183, i64 %171, i64 2, !dbg !29
+  %185 = insertelement <8 x i64> %184, i64 %173, i64 3, !dbg !29
+  %186 = insertelement <8 x i64> %185, i64 %175, i64 4, !dbg !29
+  %187 = insertelement <8 x i64> %186, i64 %177, i64 5, !dbg !29
+  %188 = insertelement <8 x i64> %187, i64 %179, i64 6, !dbg !29
+  %189 = insertelement <8 x i64> %188, i64 %181, i64 7, !dbg !29
+  %190 = icmp slt <8 x i64> %189, zeroinitializer, !dbg !29
+  %191 = insertelement <8 x i64> poison, i64 %7, i64 0, !dbg !30
+  %192 = shufflevector <8 x i64> %191, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !30
+  %193 = select <8 x i1> %190, <8 x i64> %192, <8 x i64> zeroinitializer, !dbg !30
+  %194 = add <8 x i64> %193, %189, !dbg !30
+  %195 = icmp slt <8 x i64> %194, zeroinitializer, !dbg !31
+  %196 = icmp sge <8 x i64> %194, %192, !dbg !32
+  %197 = or <8 x i1> %195, %196, !dbg !33
+  %198 = and <8 x i1> %142, %197, !dbg !34
+  %199 = bitcast <8 x i1> %198 to i8, !dbg !35
+  %.not = icmp eq i8 %199, 0, !dbg !35
+  br i1 %.not, label %201, label %200, !dbg !35
+
+200:                                              ; preds = %12
+  tail call void @__assertfail(ptr nonnull @assertMessage_0, ptr nonnull @assertFile_0, i32 36, ptr nonnull @assertFunc_0, i64 1), !dbg !35
+  unreachable, !dbg !35
+
+201:                                              ; preds = %12
+  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !35
+  %202 = extractelement <8 x i64> %40, i64 0, !dbg !36
+  %203 = sub nsw i64 %202, %121, !dbg !36
+  %204 = extractelement <8 x i64> %40, i64 1, !dbg !36
+  %205 = sub nsw i64 %204, %121, !dbg !36
+  %206 = extractelement <8 x i64> %40, i64 2, !dbg !36
+  %207 = sub nsw i64 %206, %121, !dbg !36
+  %208 = extractelement <8 x i64> %40, i64 3, !dbg !36
+  %209 = sub nsw i64 %208, %121, !dbg !36
+  %210 = extractelement <8 x i64> %40, i64 4, !dbg !36
+  %211 = sub nsw i64 %210, %121, !dbg !36
+  %212 = extractelement <8 x i64> %40, i64 5, !dbg !36
+  %213 = sub nsw i64 %212, %121, !dbg !36
+  %214 = extractelement <8 x i64> %40, i64 6, !dbg !36
+  %215 = sub nsw i64 %214, %121, !dbg !36
+  %216 = extractelement <8 x i64> %40, i64 7, !dbg !36
+  %217 = sub nsw i64 %216, %121, !dbg !36
+  %218 = extractelement <8 x i64> %194, i64 0, !dbg !37
+  %219 = mul i64 %218, %5, !dbg !37
+  %220 = extractelement <8 x i64> %194, i64 1, !dbg !37
+  %221 = mul i64 %220, %5, !dbg !37
+  %222 = extractelement <8 x i64> %194, i64 2, !dbg !37
+  %223 = mul i64 %222, %5, !dbg !37
+  %224 = extractelement <8 x i64> %194, i64 3, !dbg !37
+  %225 = mul i64 %224, %5, !dbg !37
+  %226 = extractelement <8 x i64> %194, i64 4, !dbg !37
+  %227 = mul i64 %226, %5, !dbg !37
+  %228 = extractelement <8 x i64> %194, i64 5, !dbg !37
+  %229 = mul i64 %228, %5, !dbg !37
+  %230 = extractelement <8 x i64> %194, i64 6, !dbg !37
+  %231 = mul i64 %230, %5, !dbg !37
+  %232 = extractelement <8 x i64> %194, i64 7, !dbg !37
+  %233 = mul i64 %232, %5, !dbg !37
+  %234 = getelementptr bfloat, ptr addrspace(1) %2, i64 %203, !dbg !38
+  %235 = getelementptr bfloat, ptr addrspace(1) %234, i64 %219, !dbg !38
+  %236 = getelementptr bfloat, ptr addrspace(1) %2, i64 %205, !dbg !38
+  %237 = getelementptr bfloat, ptr addrspace(1) %236, i64 %221, !dbg !38
+  %238 = getelementptr bfloat, ptr addrspace(1) %2, i64 %207, !dbg !38
+  %239 = getelementptr bfloat, ptr addrspace(1) %238, i64 %223, !dbg !38
+  %240 = getelementptr bfloat, ptr addrspace(1) %2, i64 %209, !dbg !38
+  %241 = getelementptr bfloat, ptr addrspace(1) %240, i64 %225, !dbg !38
+  %242 = getelementptr bfloat, ptr addrspace(1) %2, i64 %211, !dbg !38
+  %243 = getelementptr bfloat, ptr addrspace(1) %242, i64 %227, !dbg !38
+  %244 = getelementptr bfloat, ptr addrspace(1) %2, i64 %213, !dbg !38
+  %245 = getelementptr bfloat, ptr addrspace(1) %244, i64 %229, !dbg !38
+  %246 = getelementptr bfloat, ptr addrspace(1) %2, i64 %215, !dbg !38
+  %247 = getelementptr bfloat, ptr addrspace(1) %246, i64 %231, !dbg !38
+  %248 = getelementptr bfloat, ptr addrspace(1) %2, i64 %217, !dbg !38
+  %249 = getelementptr bfloat, ptr addrspace(1) %248, i64 %233, !dbg !38
+  %250 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %251 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %235, i64 %250, i1 %143) #4, !dbg !39
+  %252 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %253 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %237, i64 %252, i1 %146) #4, !dbg !39
+  %254 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %255 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %239, i64 %254, i1 %149) #4, !dbg !39
+  %256 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %257 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %241, i64 %256, i1 %152) #4, !dbg !39
+  %258 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %259 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %243, i64 %258, i1 %155) #4, !dbg !39
+  %260 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %261 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %245, i64 %260, i1 %158) #4, !dbg !39
+  %262 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %263 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %247, i64 %262, i1 %161) #4, !dbg !39
+  %264 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !39
+  %265 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %249, i64 %264, i1 %164) #4, !dbg !39
+  %266 = icmp slt <8 x i64> %40, %140, !dbg !40
+  %267 = add i64 %5, %41, !dbg !41
+  %268 = add i64 %5, %43, !dbg !41
+  %269 = add i64 %5, %45, !dbg !41
+  %270 = add i64 %5, %47, !dbg !41
+  %271 = add i64 %5, %49, !dbg !41
+  %272 = add i64 %5, %51, !dbg !41
+  %273 = add i64 %5, %53, !dbg !41
+  %274 = add i64 %5, %55, !dbg !41
+  %275 = sub i64 %267, %121, !dbg !42
+  %276 = sub i64 %268, %121, !dbg !42
+  %277 = sub i64 %269, %121, !dbg !42
+  %278 = sub i64 %270, %121, !dbg !42
+  %279 = sub i64 %271, %121, !dbg !42
+  %280 = sub i64 %272, %121, !dbg !42
+  %281 = sub i64 %273, %121, !dbg !42
+  %282 = sub i64 %274, %121, !dbg !42
+  %283 = getelementptr bfloat, ptr addrspace(1) %0, i64 %275, !dbg !43
+  %284 = getelementptr bfloat, ptr addrspace(1) %0, i64 %276, !dbg !43
+  %285 = getelementptr bfloat, ptr addrspace(1) %0, i64 %277, !dbg !43
+  %286 = getelementptr bfloat, ptr addrspace(1) %0, i64 %278, !dbg !43
+  %287 = getelementptr bfloat, ptr addrspace(1) %0, i64 %279, !dbg !43
+  %288 = getelementptr bfloat, ptr addrspace(1) %0, i64 %280, !dbg !43
+  %289 = getelementptr bfloat, ptr addrspace(1) %0, i64 %281, !dbg !43
+  %290 = getelementptr bfloat, ptr addrspace(1) %0, i64 %282, !dbg !43
+  %291 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %292 = and <8 x i1> %36, %266, !dbg !45
+  %293 = extractelement <8 x i1> %292, i64 0, !dbg !46
+  %294 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %283, i64 %291, i1 %293) #4, !dbg !44
+  %295 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %296 = extractelement <8 x i1> %292, i64 1, !dbg !46
+  %297 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %284, i64 %295, i1 %296) #4, !dbg !44
+  %298 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %299 = extractelement <8 x i1> %292, i64 2, !dbg !46
+  %300 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %285, i64 %298, i1 %299) #4, !dbg !44
+  %301 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %302 = extractelement <8 x i1> %292, i64 3, !dbg !46
+  %303 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %286, i64 %301, i1 %302) #4, !dbg !44
+  %304 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %305 = extractelement <8 x i1> %292, i64 4, !dbg !46
+  %306 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %287, i64 %304, i1 %305) #4, !dbg !44
+  %307 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %308 = extractelement <8 x i1> %292, i64 5, !dbg !46
+  %309 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %288, i64 %307, i1 %308) #4, !dbg !44
+  %310 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %311 = extractelement <8 x i1> %292, i64 6, !dbg !46
+  %312 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %289, i64 %310, i1 %311) #4, !dbg !44
+  %313 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !44
+  %314 = extractelement <8 x i1> %292, i64 7, !dbg !46
+  %315 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %290, i64 %313, i1 %314) #4, !dbg !44
+  %316 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %317 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %97, i64 %316, i1 %293) #4, !dbg !46
+  %318 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %319 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %98, i64 %318, i1 %296) #4, !dbg !46
+  %320 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %321 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %99, i64 %320, i1 %299) #4, !dbg !46
+  %322 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %323 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %100, i64 %322, i1 %302) #4, !dbg !46
+  %324 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %325 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %101, i64 %324, i1 %305) #4, !dbg !46
+  %326 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %327 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %102, i64 %326, i1 %308) #4, !dbg !46
+  %328 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %329 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %103, i64 %328, i1 %311) #4, !dbg !46
+  %330 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !46
+  %331 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %104, i64 %330, i1 %314) #4, !dbg !46
+  %332 = insertelement <8 x i64> poison, i64 %317, i64 0, !dbg !47
+  %333 = insertelement <8 x i64> %332, i64 %319, i64 1, !dbg !47
+  %334 = insertelement <8 x i64> %333, i64 %321, i64 2, !dbg !47
+  %335 = insertelement <8 x i64> %334, i64 %323, i64 3, !dbg !47
+  %336 = insertelement <8 x i64> %335, i64 %325, i64 4, !dbg !47
+  %337 = insertelement <8 x i64> %336, i64 %327, i64 5, !dbg !47
+  %338 = insertelement <8 x i64> %337, i64 %329, i64 6, !dbg !47
+  %339 = insertelement <8 x i64> %338, i64 %331, i64 7, !dbg !47
+  %340 = icmp slt <8 x i64> %339, zeroinitializer, !dbg !47
+  %341 = select <8 x i1> %340, <8 x i64> %192, <8 x i64> zeroinitializer, !dbg !48
+  %342 = add <8 x i64> %341, %339, !dbg !48
+  %343 = icmp slt <8 x i64> %342, zeroinitializer, !dbg !49
+  %344 = icmp sge <8 x i64> %342, %192, !dbg !50
+  %345 = or <8 x i1> %343, %344, !dbg !51
+  %346 = and <8 x i1> %292, %345, !dbg !52
+  %347 = bitcast <8 x i1> %346 to i8, !dbg !53
+  %.not97 = icmp eq i8 %347, 0, !dbg !53
+  br i1 %.not97, label %349, label %348, !dbg !53
+
+348:                                              ; preds = %201
+  tail call void @__assertfail(ptr nonnull @assertMessage_1, ptr nonnull @assertFile_1, i32 51, ptr nonnull @assertFunc_1, i64 1), !dbg !53
+  unreachable, !dbg !53
+
+349:                                              ; preds = %201
+  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !53
+  %350 = sub i64 %5, %121, !dbg !54
+  %351 = extractelement <8 x i64> %342, i64 0, !dbg !55
+  %352 = mul i64 %351, %5, !dbg !55
+  %353 = extractelement <8 x i64> %342, i64 1, !dbg !55
+  %354 = mul i64 %353, %5, !dbg !55
+  %355 = extractelement <8 x i64> %342, i64 2, !dbg !55
+  %356 = mul i64 %355, %5, !dbg !55
+  %357 = extractelement <8 x i64> %342, i64 3, !dbg !55
+  %358 = mul i64 %357, %5, !dbg !55
+  %359 = extractelement <8 x i64> %342, i64 4, !dbg !55
+  %360 = mul i64 %359, %5, !dbg !55
+  %361 = extractelement <8 x i64> %342, i64 5, !dbg !55
+  %362 = mul i64 %361, %5, !dbg !55
+  %363 = extractelement <8 x i64> %342, i64 6, !dbg !55
+  %364 = mul i64 %363, %5, !dbg !55
+  %365 = extractelement <8 x i64> %342, i64 7, !dbg !55
+  %366 = mul i64 %365, %5, !dbg !55
+  %367 = getelementptr bfloat, ptr addrspace(1) %2, i64 %350, !dbg !56
+  %368 = getelementptr bfloat, ptr addrspace(1) %367, i64 %202, !dbg !56
+  %369 = getelementptr bfloat, ptr addrspace(1) %368, i64 %352, !dbg !56
+  %370 = getelementptr bfloat, ptr addrspace(1) %367, i64 %204, !dbg !56
+  %371 = getelementptr bfloat, ptr addrspace(1) %370, i64 %354, !dbg !56
+  %372 = getelementptr bfloat, ptr addrspace(1) %367, i64 %206, !dbg !56
+  %373 = getelementptr bfloat, ptr addrspace(1) %372, i64 %356, !dbg !56
+  %374 = getelementptr bfloat, ptr addrspace(1) %367, i64 %208, !dbg !56
+  %375 = getelementptr bfloat, ptr addrspace(1) %374, i64 %358, !dbg !56
+  %376 = getelementptr bfloat, ptr addrspace(1) %367, i64 %210, !dbg !56
+  %377 = getelementptr bfloat, ptr addrspace(1) %376, i64 %360, !dbg !56
+  %378 = getelementptr bfloat, ptr addrspace(1) %367, i64 %212, !dbg !56
+  %379 = getelementptr bfloat, ptr addrspace(1) %378, i64 %362, !dbg !56
+  %380 = getelementptr bfloat, ptr addrspace(1) %367, i64 %214, !dbg !56
+  %381 = getelementptr bfloat, ptr addrspace(1) %380, i64 %364, !dbg !56
+  %382 = getelementptr bfloat, ptr addrspace(1) %367, i64 %216, !dbg !56
+  %383 = getelementptr bfloat, ptr addrspace(1) %382, i64 %366, !dbg !56
+  %384 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %385 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %369, i64 %384, i1 %293) #4, !dbg !57
+  %386 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %387 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %371, i64 %386, i1 %296) #4, !dbg !57
+  %388 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %389 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %373, i64 %388, i1 %299) #4, !dbg !57
+  %390 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %391 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %375, i64 %390, i1 %302) #4, !dbg !57
+  %392 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %393 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %377, i64 %392, i1 %305) #4, !dbg !57
+  %394 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %395 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %379, i64 %394, i1 %308) #4, !dbg !57
+  %396 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %397 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %381, i64 %396, i1 %311) #4, !dbg !57
+  %398 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
+  %399 = tail call i16 asm sideeffect "mov.u16 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $2 + 0 ], $3;", "=c,c,l,l,b"(i16 0, ptr addrspace(1) %383, i64 %398, i1 %314) #4, !dbg !57
+  %400 = insertelement <8 x i64> poison, i64 %106, i64 0, !dbg !58
+  %401 = insertelement <8 x i64> %400, i64 %108, i64 1, !dbg !58
+  %402 = insertelement <8 x i64> %401, i64 %110, i64 2, !dbg !58
+  %403 = insertelement <8 x i64> %402, i64 %112, i64 3, !dbg !58
+  %404 = insertelement <8 x i64> %403, i64 %114, i64 4, !dbg !58
+  %405 = insertelement <8 x i64> %404, i64 %116, i64 5, !dbg !58
+  %406 = insertelement <8 x i64> %405, i64 %118, i64 6, !dbg !58
+  %407 = insertelement <8 x i64> %406, i64 %120, i64 7, !dbg !58
+  %408 = icmp slt <8 x i64> %407, zeroinitializer, !dbg !58
+  %409 = insertelement <8 x i64> poison, i64 %8, i64 0, !dbg !59
+  %410 = shufflevector <8 x i64> %409, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !59
+  %411 = select <8 x i1> %408, <8 x i64> %410, <8 x i64> zeroinitializer, !dbg !59
+  %412 = add <8 x i64> %411, %407, !dbg !59
+  %413 = icmp slt <8 x i64> %412, zeroinitializer, !dbg !60
+  %414 = icmp sge <8 x i64> %412, %410, !dbg !61
+  %415 = or <8 x i1> %413, %414, !dbg !62
+  %416 = and <8 x i1> %36, %415, !dbg !63
+  %417 = bitcast <8 x i1> %416 to i8, !dbg !64
+  %.not98 = icmp eq i8 %417, 0, !dbg !64
+  br i1 %.not98, label %419, label %418, !dbg !64
+
+418:                                              ; preds = %349
+  tail call void @__assertfail(ptr nonnull @assertMessage_2, ptr nonnull @assertFile_2, i32 62, ptr nonnull @assertFunc_2, i64 1), !dbg !64
+  unreachable, !dbg !64
+
+419:                                              ; preds = %349
+  %420 = bitcast i16 %165 to bfloat, !dbg !25
+  %421 = fpext bfloat %420 to float, !dbg !65
+  %422 = bitcast i16 %265 to bfloat, !dbg !39
+  %423 = fpext bfloat %422 to float, !dbg !66
+  %424 = fmul float %421, %423, !dbg !67
+  %425 = fsub float 0.000000e+00, %424, !dbg !68
+  %426 = extractelement <8 x i1> %141, i64 7, !dbg !69
+  %427 = select i1 %426, float %425, float 0.000000e+00, !dbg !69
+  %428 = bitcast i16 %315 to bfloat, !dbg !44
+  %429 = fpext bfloat %428 to float, !dbg !70
+  %430 = bitcast i16 %399 to bfloat, !dbg !57
+  %431 = fpext bfloat %430 to float, !dbg !71
+  %432 = fmul float %429, %431, !dbg !72
+  %433 = extractelement <8 x i1> %266, i64 7, !dbg !69
+  %434 = select i1 %433, float %432, float 0.000000e+00, !dbg !69
+  %435 = fadd float %427, %434, !dbg !73
+  %436 = bitcast i16 %162 to bfloat, !dbg !25
+  %437 = fpext bfloat %436 to float, !dbg !65
+  %438 = bitcast i16 %263 to bfloat, !dbg !39
+  %439 = fpext bfloat %438 to float, !dbg !66
+  %440 = fmul float %437, %439, !dbg !67
+  %441 = fsub float 0.000000e+00, %440, !dbg !68
+  %442 = extractelement <8 x i1> %141, i64 6, !dbg !69
+  %443 = select i1 %442, float %441, float 0.000000e+00, !dbg !69
+  %444 = bitcast i16 %312 to bfloat, !dbg !44
+  %445 = fpext bfloat %444 to float, !dbg !70
+  %446 = bitcast i16 %397 to bfloat, !dbg !57
+  %447 = fpext bfloat %446 to float, !dbg !71
+  %448 = fmul float %445, %447, !dbg !72
+  %449 = extractelement <8 x i1> %266, i64 6, !dbg !69
+  %450 = select i1 %449, float %448, float 0.000000e+00, !dbg !69
+  %451 = fadd float %443, %450, !dbg !73
+  %452 = bitcast i16 %159 to bfloat, !dbg !25
+  %453 = fpext bfloat %452 to float, !dbg !65
+  %454 = bitcast i16 %261 to bfloat, !dbg !39
+  %455 = fpext bfloat %454 to float, !dbg !66
+  %456 = fmul float %453, %455, !dbg !67
+  %457 = fsub float 0.000000e+00, %456, !dbg !68
+  %458 = extractelement <8 x i1> %141, i64 5, !dbg !69
+  %459 = select i1 %458, float %457, float 0.000000e+00, !dbg !69
+  %460 = bitcast i16 %309 to bfloat, !dbg !44
+  %461 = fpext bfloat %460 to float, !dbg !70
+  %462 = bitcast i16 %395 to bfloat, !dbg !57
+  %463 = fpext bfloat %462 to float, !dbg !71
+  %464 = fmul float %461, %463, !dbg !72
+  %465 = extractelement <8 x i1> %266, i64 5, !dbg !69
+  %466 = select i1 %465, float %464, float 0.000000e+00, !dbg !69
+  %467 = fadd float %459, %466, !dbg !73
+  %468 = bitcast i16 %156 to bfloat, !dbg !25
+  %469 = fpext bfloat %468 to float, !dbg !65
+  %470 = bitcast i16 %259 to bfloat, !dbg !39
+  %471 = fpext bfloat %470 to float, !dbg !66
+  %472 = fmul float %469, %471, !dbg !67
+  %473 = fsub float 0.000000e+00, %472, !dbg !68
+  %474 = extractelement <8 x i1> %141, i64 4, !dbg !69
+  %475 = select i1 %474, float %473, float 0.000000e+00, !dbg !69
+  %476 = bitcast i16 %306 to bfloat, !dbg !44
+  %477 = fpext bfloat %476 to float, !dbg !70
+  %478 = bitcast i16 %393 to bfloat, !dbg !57
+  %479 = fpext bfloat %478 to float, !dbg !71
+  %480 = fmul float %477, %479, !dbg !72
+  %481 = extractelement <8 x i1> %266, i64 4, !dbg !69
+  %482 = select i1 %481, float %480, float 0.000000e+00, !dbg !69
+  %483 = fadd float %475, %482, !dbg !73
+  %484 = bitcast i16 %153 to bfloat, !dbg !25
+  %485 = fpext bfloat %484 to float, !dbg !65
+  %486 = bitcast i16 %257 to bfloat, !dbg !39
+  %487 = fpext bfloat %486 to float, !dbg !66
+  %488 = fmul float %485, %487, !dbg !67
+  %489 = fsub float 0.000000e+00, %488, !dbg !68
+  %490 = extractelement <8 x i1> %141, i64 3, !dbg !69
+  %491 = select i1 %490, float %489, float 0.000000e+00, !dbg !69
+  %492 = bitcast i16 %303 to bfloat, !dbg !44
+  %493 = fpext bfloat %492 to float, !dbg !70
+  %494 = bitcast i16 %391 to bfloat, !dbg !57
+  %495 = fpext bfloat %494 to float, !dbg !71
+  %496 = fmul float %493, %495, !dbg !72
+  %497 = extractelement <8 x i1> %266, i64 3, !dbg !69
+  %498 = select i1 %497, float %496, float 0.000000e+00, !dbg !69
+  %499 = fadd float %491, %498, !dbg !73
+  %500 = bitcast i16 %150 to bfloat, !dbg !25
+  %501 = fpext bfloat %500 to float, !dbg !65
+  %502 = bitcast i16 %255 to bfloat, !dbg !39
+  %503 = fpext bfloat %502 to float, !dbg !66
+  %504 = fmul float %501, %503, !dbg !67
+  %505 = fsub float 0.000000e+00, %504, !dbg !68
+  %506 = extractelement <8 x i1> %141, i64 2, !dbg !69
+  %507 = select i1 %506, float %505, float 0.000000e+00, !dbg !69
+  %508 = bitcast i16 %300 to bfloat, !dbg !44
+  %509 = fpext bfloat %508 to float, !dbg !70
+  %510 = bitcast i16 %389 to bfloat, !dbg !57
+  %511 = fpext bfloat %510 to float, !dbg !71
+  %512 = fmul float %509, %511, !dbg !72
+  %513 = extractelement <8 x i1> %266, i64 2, !dbg !69
+  %514 = select i1 %513, float %512, float 0.000000e+00, !dbg !69
+  %515 = fadd float %507, %514, !dbg !73
+  %516 = bitcast i16 %147 to bfloat, !dbg !25
+  %517 = fpext bfloat %516 to float, !dbg !65
+  %518 = bitcast i16 %253 to bfloat, !dbg !39
+  %519 = fpext bfloat %518 to float, !dbg !66
+  %520 = fmul float %517, %519, !dbg !67
+  %521 = fsub float 0.000000e+00, %520, !dbg !68
+  %522 = extractelement <8 x i1> %141, i64 1, !dbg !69
+  %523 = select i1 %522, float %521, float 0.000000e+00, !dbg !69
+  %524 = bitcast i16 %297 to bfloat, !dbg !44
+  %525 = fpext bfloat %524 to float, !dbg !70
+  %526 = bitcast i16 %387 to bfloat, !dbg !57
+  %527 = fpext bfloat %526 to float, !dbg !71
+  %528 = fmul float %525, %527, !dbg !72
+  %529 = extractelement <8 x i1> %266, i64 1, !dbg !69
+  %530 = select i1 %529, float %528, float 0.000000e+00, !dbg !69
+  %531 = fadd float %523, %530, !dbg !73
+  %532 = bitcast i16 %144 to bfloat, !dbg !25
+  %533 = fpext bfloat %532 to float, !dbg !65
+  %534 = bitcast i16 %251 to bfloat, !dbg !39
+  %535 = fpext bfloat %534 to float, !dbg !66
+  %536 = fmul float %533, %535, !dbg !67
+  %537 = fsub float 0.000000e+00, %536, !dbg !68
+  %538 = extractelement <8 x i1> %141, i64 0, !dbg !69
+  %539 = select i1 %538, float %537, float 0.000000e+00, !dbg !69
+  %540 = bitcast i16 %294 to bfloat, !dbg !44
+  %541 = fpext bfloat %540 to float, !dbg !70
+  %542 = bitcast i16 %385 to bfloat, !dbg !57
+  %543 = fpext bfloat %542 to float, !dbg !71
+  %544 = fmul float %541, %543, !dbg !72
+  %545 = extractelement <8 x i1> %266, i64 0, !dbg !69
+  %546 = select i1 %545, float %544, float 0.000000e+00, !dbg !69
+  %547 = fadd float %539, %546, !dbg !73
+  %548 = bitcast i16 %96 to bfloat, !dbg !19
+  %549 = fpext bfloat %548 to float, !dbg !74
+  %550 = bitcast i16 %93 to bfloat, !dbg !19
+  %551 = fpext bfloat %550 to float, !dbg !74
+  %552 = bitcast i16 %90 to bfloat, !dbg !19
+  %553 = fpext bfloat %552 to float, !dbg !74
+  %554 = bitcast i16 %87 to bfloat, !dbg !19
+  %555 = fpext bfloat %554 to float, !dbg !74
+  %556 = bitcast i16 %84 to bfloat, !dbg !19
+  %557 = fpext bfloat %556 to float, !dbg !74
+  %558 = bitcast i16 %81 to bfloat, !dbg !19
+  %559 = fpext bfloat %558 to float, !dbg !74
+  %560 = bitcast i16 %78 to bfloat, !dbg !19
+  %561 = fpext bfloat %560 to float, !dbg !74
+  %562 = bitcast i16 %75 to bfloat, !dbg !19
+  %563 = fpext bfloat %562 to float, !dbg !74
+  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !64
+  %564 = extractelement <8 x i64> %412, i64 0, !dbg !75
+  %565 = mul i64 %564, %5, !dbg !75
+  %566 = extractelement <8 x i64> %412, i64 1, !dbg !75
+  %567 = mul i64 %566, %5, !dbg !75
+  %568 = extractelement <8 x i64> %412, i64 2, !dbg !75
+  %569 = mul i64 %568, %5, !dbg !75
+  %570 = extractelement <8 x i64> %412, i64 3, !dbg !75
+  %571 = mul i64 %570, %5, !dbg !75
+  %572 = extractelement <8 x i64> %412, i64 4, !dbg !75
+  %573 = mul i64 %572, %5, !dbg !75
+  %574 = extractelement <8 x i64> %412, i64 5, !dbg !75
+  %575 = mul i64 %574, %5, !dbg !75
+  %576 = extractelement <8 x i64> %412, i64 6, !dbg !75
+  %577 = mul i64 %576, %5, !dbg !75
+  %578 = extractelement <8 x i64> %412, i64 7, !dbg !75
+  %579 = mul i64 %578, %5, !dbg !75
+  %580 = getelementptr bfloat, ptr addrspace(1) %3, i64 %202, !dbg !76
+  %581 = getelementptr bfloat, ptr addrspace(1) %580, i64 %565, !dbg !76
+  %582 = getelementptr bfloat, ptr addrspace(1) %3, i64 %204, !dbg !76
+  %583 = getelementptr bfloat, ptr addrspace(1) %582, i64 %567, !dbg !76
+  %584 = getelementptr bfloat, ptr addrspace(1) %3, i64 %206, !dbg !76
+  %585 = getelementptr bfloat, ptr addrspace(1) %584, i64 %569, !dbg !76
+  %586 = getelementptr bfloat, ptr addrspace(1) %3, i64 %208, !dbg !76
+  %587 = getelementptr bfloat, ptr addrspace(1) %586, i64 %571, !dbg !76
+  %588 = getelementptr bfloat, ptr addrspace(1) %3, i64 %210, !dbg !76
+  %589 = getelementptr bfloat, ptr addrspace(1) %588, i64 %573, !dbg !76
+  %590 = getelementptr bfloat, ptr addrspace(1) %3, i64 %212, !dbg !76
+  %591 = getelementptr bfloat, ptr addrspace(1) %590, i64 %575, !dbg !76
+  %592 = getelementptr bfloat, ptr addrspace(1) %3, i64 %214, !dbg !76
+  %593 = getelementptr bfloat, ptr addrspace(1) %592, i64 %577, !dbg !76
+  %594 = getelementptr bfloat, ptr addrspace(1) %3, i64 %216, !dbg !76
+  %595 = getelementptr bfloat, ptr addrspace(1) %594, i64 %579, !dbg !76
+  %596 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %597 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %581, i64 %596, i1 %74) #4, !dbg !77
+  %598 = bitcast i16 %597 to bfloat, !dbg !77
+  %599 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %600 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %583, i64 %599, i1 %77) #4, !dbg !77
+  %601 = bitcast i16 %600 to bfloat, !dbg !77
+  %602 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %603 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %585, i64 %602, i1 %80) #4, !dbg !77
+  %604 = bitcast i16 %603 to bfloat, !dbg !77
+  %605 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %606 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %587, i64 %605, i1 %83) #4, !dbg !77
+  %607 = bitcast i16 %606 to bfloat, !dbg !77
+  %608 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %609 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %589, i64 %608, i1 %86) #4, !dbg !77
+  %610 = bitcast i16 %609 to bfloat, !dbg !77
+  %611 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %612 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %591, i64 %611, i1 %89) #4, !dbg !77
+  %613 = bitcast i16 %612 to bfloat, !dbg !77
+  %614 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %615 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %593, i64 %614, i1 %92) #4, !dbg !77
+  %616 = bitcast i16 %615 to bfloat, !dbg !77
+  %617 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !77
+  %618 = tail call i16 asm sideeffect "mov.u16 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b16 { $0 }, [ $1 + 0 ], $2;", "=c,l,l,b"(ptr addrspace(1) %595, i64 %617, i1 %95) #4, !dbg !77
+  %619 = bitcast i16 %618 to bfloat, !dbg !77
+  %620 = fpext bfloat %598 to float, !dbg !78
+  %621 = fpext bfloat %601 to float, !dbg !78
+  %622 = fpext bfloat %604 to float, !dbg !78
+  %623 = fpext bfloat %607 to float, !dbg !78
+  %624 = fpext bfloat %610 to float, !dbg !78
+  %625 = fpext bfloat %613 to float, !dbg !78
+  %626 = fpext bfloat %616 to float, !dbg !78
+  %627 = fpext bfloat %619 to float, !dbg !78
+  %628 = fmul float %563, %620, !dbg !79
+  %629 = fmul float %561, %621, !dbg !79
+  %630 = fmul float %559, %622, !dbg !79
+  %631 = fmul float %557, %623, !dbg !79
+  %632 = fmul float %555, %624, !dbg !79
+  %633 = fmul float %553, %625, !dbg !79
+  %634 = fmul float %551, %626, !dbg !79
+  %635 = fmul float %549, %627, !dbg !79
+  %636 = fadd float %547, %628, !dbg !80
+  %637 = fadd float %531, %629, !dbg !80
+  %638 = fadd float %515, %630, !dbg !80
+  %639 = fadd float %499, %631, !dbg !80
+  %640 = fadd float %483, %632, !dbg !80
+  %641 = fadd float %467, %633, !dbg !80
+  %642 = fadd float %451, %634, !dbg !80
+  %643 = fadd float %435, %635, !dbg !80
+  %644 = getelementptr bfloat, ptr addrspace(1) %4, i64 %41, !dbg !81
+  %645 = getelementptr bfloat, ptr addrspace(1) %4, i64 %43, !dbg !81
+  %646 = getelementptr bfloat, ptr addrspace(1) %4, i64 %45, !dbg !81
+  %647 = getelementptr bfloat, ptr addrspace(1) %4, i64 %47, !dbg !81
+  %648 = getelementptr bfloat, ptr addrspace(1) %4, i64 %49, !dbg !81
+  %649 = getelementptr bfloat, ptr addrspace(1) %4, i64 %51, !dbg !81
+  %650 = getelementptr bfloat, ptr addrspace(1) %4, i64 %53, !dbg !81
+  %651 = getelementptr bfloat, ptr addrspace(1) %4, i64 %55, !dbg !81
+  %652 = fptrunc float %636 to bfloat, !dbg !82
+  %653 = fptrunc float %637 to bfloat, !dbg !82
+  %654 = fptrunc float %638 to bfloat, !dbg !82
+  %655 = fptrunc float %639 to bfloat, !dbg !82
+  %656 = fptrunc float %640 to bfloat, !dbg !82
+  %657 = fptrunc float %641 to bfloat, !dbg !82
+  %658 = fptrunc float %642 to bfloat, !dbg !82
+  %659 = fptrunc float %643 to bfloat, !dbg !82
+  %660 = bitcast bfloat %652 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %660, ptr addrspace(1) %644, i1 %74) #4, !dbg !82
+  %661 = bitcast bfloat %653 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %661, ptr addrspace(1) %645, i1 %77) #4, !dbg !82
+  %662 = bitcast bfloat %654 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %662, ptr addrspace(1) %646, i1 %80) #4, !dbg !82
+  %663 = bitcast bfloat %655 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %663, ptr addrspace(1) %647, i1 %83) #4, !dbg !82
+  %664 = bitcast bfloat %656 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %664, ptr addrspace(1) %648, i1 %86) #4, !dbg !82
+  %665 = bitcast bfloat %657 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %665, ptr addrspace(1) %649, i1 %89) #4, !dbg !82
+  %666 = bitcast bfloat %658 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %666, ptr addrspace(1) %650, i1 %92) #4, !dbg !82
+  %667 = bitcast bfloat %659 to i16, !dbg !82
+  tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %667, ptr addrspace(1) %651, i1 %95) #4, !dbg !82
+  ret void, !dbg !83
+}
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #2
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #2
+
+; Function Attrs: convergent nocallback nounwind
+declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3
+
+attributes #0 = { noreturn }
+attributes #1 = { "nvvm.reqntid"="128" }
+attributes #2 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+attributes #3 = { convergent nocallback nounwind }
+attributes #4 = { nounwind }
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+!llvm.ident = !{!4}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
+!1 = !DIFile(filename: "cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb")
+!2 = !{i32 2, !"Debug Info Version", i32 3}
+!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
+!4 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}
+!5 = !DISubprogram(name: "__assertfail", linkageName: "__assertfail", scope: !6, file: !6, type: !7, spFlags: DISPFlagOptimized)
+!6 = !DIFile(filename: "<unknown>", directory: "")
+!7 = !DISubroutineType(cc: DW_CC_normal, types: !8)
+!8 = !{}
+!9 = distinct !DISubprogram(name: "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0", linkageName: "triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0", scope: !1, file: !1, line: 18, type: !7, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
+!10 = !DILocation(line: 19, column: 28, scope: !9)
+!11 = !DILocation(line: 19, column: 33, scope: !9)
+!12 = !DILocation(line: 20, column: 36, scope: !9)
+!13 = !DILocation(line: 21, column: 21, scope: !9)
+!14 = !DILocation(line: 20, column: 23, scope: !9)
+!15 = !DILocation(line: 22, column: 19, scope: !9)
+!16 = !DILocation(line: 25, column: 31, scope: !9)
+!17 = !DILocation(line: 24, column: 21, scope: !9)
+!18 = !DILocation(line: 24, column: 28, scope: !9)
+!19 = !DILocation(line: 25, column: 36, scope: !9)
+!20 = !DILocation(line: 26, column: 31, scope: !9)
+!21 = !DILocation(line: 26, column: 36, scope: !9)
+!22 = !DILocation(line: 28, column: 18, scope: !9)
+!23 = !DILocation(line: 30, column: 35, scope: !9)
+!24 = !DILocation(line: 30, column: 30, scope: !9)
+!25 = !DILocation(line: 30, column: 53, scope: !9)
+!26 = !DILocation(line: 29, column: 19, scope: !9)
+!27 = !DILocation(line: 30, column: 60, scope: !9)
+!28 = !DILocation(line: 31, column: 35, scope: !9)
+!29 = !DILocation(line: 34, column: 18, scope: !9)
+!30 = !DILocation(line: 35, column: 32, scope: !9)
+!31 = !DILocation(line: 36, column: 28, scope: !9)
+!32 = !DILocation(line: 36, column: 98, scope: !9)
+!33 = !DILocation(line: 36, column: 64, scope: !9)
+!34 = !DILocation(line: 36, column: 106, scope: !9)
+!35 = !DILocation(line: 36, column: 123, scope: !9)
+!36 = !DILocation(line: 37, column: 36, scope: !9)
+!37 = !DILocation(line: 37, column: 58, scope: !9)
+!38 = !DILocation(line: 37, column: 31, scope: !9)
+!39 = !DILocation(line: 37, column: 65, scope: !9)
+!40 = !DILocation(line: 44, column: 19, scope: !9)
+!41 = !DILocation(line: 45, column: 37, scope: !9)
+!42 = !DILocation(line: 45, column: 42, scope: !9)
+!43 = !DILocation(line: 45, column: 31, scope: !9)
+!44 = !DILocation(line: 45, column: 60, scope: !9)
+!45 = !DILocation(line: 45, column: 68, scope: !9)
+!46 = !DILocation(line: 46, column: 36, scope: !9)
+!47 = !DILocation(line: 49, column: 20, scope: !9)
+!48 = !DILocation(line: 50, column: 35, scope: !9)
+!49 = !DILocation(line: 51, column: 28, scope: !9)
+!50 = !DILocation(line: 51, column: 100, scope: !9)
+!51 = !DILocation(line: 51, column: 65, scope: !9)
+!52 = !DILocation(line: 51, column: 108, scope: !9)
+!53 = !DILocation(line: 51, column: 126, scope: !9)
+!54 = !DILocation(line: 52, column: 37, scope: !9)
+!55 = !DILocation(line: 52, column: 64, scope: !9)
+!56 = !DILocation(line: 52, column: 31, scope: !9)
+!57 = !DILocation(line: 52, column: 72, scope: !9)
+!58 = !DILocation(line: 60, column: 20, scope: !9)
+!59 = !DILocation(line: 61, column: 35, scope: !9)
+!60 = !DILocation(line: 62, column: 28, scope: !9)
+!61 = !DILocation(line: 62, column: 46, scope: !9)
+!62 = !DILocation(line: 62, column: 38, scope: !9)
+!63 = !DILocation(line: 62, column: 54, scope: !9)
+!64 = !DILocation(line: 62, column: 64, scope: !9)
+!65 = !DILocation(line: 30, column: 111, scope: !9)
+!66 = !DILocation(line: 37, column: 123, scope: !9)
+!67 = !DILocation(line: 38, column: 19, scope: !9)
+!68 = !DILocation(line: 39, column: 13, scope: !9)
+!69 = !DILocation(line: 0, scope: !9)
+!70 = !DILocation(line: 45, column: 119, scope: !9)
+!71 = !DILocation(line: 52, column: 131, scope: !9)
+!72 = !DILocation(line: 53, column: 20, scope: !9)
+!73 = !DILocation(line: 57, column: 20, scope: !9)
+!74 = !DILocation(line: 25, column: 76, scope: !9)
+!75 = !DILocation(line: 63, column: 40, scope: !9)
+!76 = !DILocation(line: 63, column: 31, scope: !9)
+!77 = !DILocation(line: 63, column: 48, scope: !9)
+!78 = !DILocation(line: 63, column: 88, scope: !9)
+!79 = !DILocation(line: 64, column: 20, scope: !9)
+!80 = !DILocation(line: 65, column: 20, scope: !9)
+!81 = !DILocation(line: 66, column: 25, scope: !9)
+!82 = !DILocation(line: 66, column: 37, scope: !9)
+!83 = !DILocation(line: 66, column: 4, scope: !9)

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ptx

@@ -0,0 +1,1936 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.7
+.target sm_90a
+.address_size 64
+
+	// .globl	triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0 // -- Begin function triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0
+.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
+)
+.noreturn;
+.global .align 1 .b8 assertFunc_2[8] = {117, 110, 107, 110, 111, 119, 110};
+.global .align 1 .b8 assertFile_2[114] = {47, 119, 111, 114, 107, 115, 112, 97, 99, 101, 47, 104, 97, 110, 114, 117, 105, 47, 83, 112, 101, 99, 70, 111, 114, 103, 101, 45, 101, 120, 116, 47, 99, 97, 99, 104, 101, 47, 99, 111, 109, 112, 105, 108, 101, 100, 95, 107, 101, 114, 110, 101, 108, 115, 47, 109, 98, 47, 99, 109, 98, 53, 122, 108, 108, 100, 107, 114, 102, 108, 101, 121, 104, 99, 107, 104, 105, 111, 109, 104, 112, 102, 115, 118, 107, 102, 107, 108, 108, 111, 116, 98, 111, 112, 97, 108, 100, 100, 119, 107, 100, 102, 104, 108, 105, 122, 118, 101, 117, 120, 46, 112, 121};
+.global .align 1 .b8 assertMessage_2[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, 54, 32, 60, 32, 107, 115, 51};
+.global .align 1 .b8 assertFunc_1[8] = {117, 110, 107, 110, 111, 119, 110};
+.global .align 1 .b8 assertFile_1[114] = {47, 119, 111, 114, 107, 115, 112, 97, 99, 101, 47, 104, 97, 110, 114, 117, 105, 47, 83, 112, 101, 99, 70, 111, 114, 103, 101, 45, 101, 120, 116, 47, 99, 97, 99, 104, 101, 47, 99, 111, 109, 112, 105, 108, 101, 100, 95, 107, 101, 114, 110, 101, 108, 115, 47, 109, 98, 47, 99, 109, 98, 53, 122, 108, 108, 100, 107, 114, 102, 108, 101, 121, 104, 99, 107, 104, 105, 111, 109, 104, 112, 102, 115, 118, 107, 102, 107, 108, 108, 111, 116, 98, 111, 112, 97, 108, 100, 100, 119, 107, 100, 102, 104, 108, 105, 122, 118, 101, 117, 120, 46, 112, 121};
+.global .align 1 .b8 assertMessage_1[65] = {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, 108, 46, 98, 114, 111, 97, 100, 99, 97, 115, 116, 95, 116, 111, 40, 116, 109, 112, 50, 51, 44, 32, 91, 88, 66, 76, 79, 67, 75, 93, 41, 32, 60, 32, 107, 115, 50};
+.global .align 1 .b8 assertFunc_0[8] = {117, 110, 107, 110, 111, 119, 110};
+.global .align 1 .b8 assertFile_0[114] = {47, 119, 111, 114, 107, 115, 112, 97, 99, 101, 47, 104, 97, 110, 114, 117, 105, 47, 83, 112, 101, 99, 70, 111, 114, 103, 101, 45, 101, 120, 116, 47, 99, 97, 99, 104, 101, 47, 99, 111, 109, 112, 105, 108, 101, 100, 95, 107, 101, 114, 110, 101, 108, 115, 47, 109, 98, 47, 99, 109, 98, 53, 122, 108, 108, 100, 107, 114, 102, 108, 101, 121, 104, 99, 107, 104, 105, 111, 109, 104, 112, 102, 115, 118, 107, 102, 107, 108, 108, 111, 116, 98, 111, 112, 97, 108, 100, 100, 119, 107, 100, 102, 104, 108, 105, 122, 118, 101, 117, 120, 46, 112, 121};
+.global .align 1 .b8 assertMessage_0[64] = {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, 108, 46, 98, 114, 111, 97, 100, 99, 97, 115, 116, 95, 116, 111, 40, 116, 109, 112, 56, 44, 32, 91, 88, 66, 76, 79, 67, 75, 93, 41, 32, 60, 32, 107, 115, 50};
+                                        // @triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0
+.visible .entry triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0(
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_0,
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_1,
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_2,
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_3,
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_4,
+	.param .u64 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_5,
+	.param .u64 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_6,
+	.param .u64 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_7,
+	.param .u64 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_8,
+	.param .u32 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_9,
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_10,
+	.param .u64 .ptr .global .align 1 triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_11
+)
+.reqntid 128
+{
+	.reg .pred 	%p<268>;
+	.reg .b16 	%rs<207>;
+	.reg .b32 	%r<174>;
+	.reg .b64 	%rd<682>;
+	.loc	1 18 0                          // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:18:0
+$L__func_begin0:
+	.loc	1 18 0                          // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:18:0
+
+// %bb.0:
+	ld.param.b64 	%rd143, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_5];
+$L__tmp0:
+	.loc	1 19 28                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:19:28
+	mov.u32 	%r17, %ctaid.x;
+	.loc	1 19 33                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:19:33
+	shl.b32 	%r18, %r17, 10;
+	.loc	1 20 36                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:20:36
+	mov.u32 	%r19, %tid.x;
+	shl.b32 	%r20, %r19, 3;
+	and.b32 	%r21, %r20, 1016;
+	.loc	1 20 23                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:20:23
+	or.b32 	%r9, %r21, %r18;
+	or.b32 	%r10, %r9, 1;
+	or.b32 	%r11, %r9, 2;
+	.loc	1 22 19                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:22:19
+	cvt.s64.s32 	%rd2, %r10;
+	cvt.s64.s32 	%rd1, %r9;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd147, %rd1, %rd143;
+	and.b64 	%rd148, %rd147, -4294967296;
+	setp.ne.b64 	%p25, %rd148, 0;
+	@%p25 bra 	$L__BB0_2;
+	bra.uni 	$L__BB0_1;
+$L__BB0_2:
+	div.s64 	%rd666, %rd1, %rd143;
+	bra.uni 	$L__BB0_3;
+$L__BB0_1:
+	cvt.u32.u64 	%r22, %rd143;
+	cvt.u32.u64 	%r23, %rd1;
+	div.u32 	%r24, %r23, %r22;
+	cvt.u64.u32 	%rd666, %r24;
+$L__BB0_3:
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	or.b32 	%r12, %r9, 3;
+	cvt.s64.s32 	%rd3, %r11;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd149, %rd2, %rd143;
+	and.b64 	%rd150, %rd149, -4294967296;
+	setp.ne.b64 	%p26, %rd150, 0;
+	@%p26 bra 	$L__BB0_5;
+	bra.uni 	$L__BB0_4;
+$L__BB0_5:
+	div.s64 	%rd667, %rd2, %rd143;
+	bra.uni 	$L__BB0_6;
+$L__BB0_4:
+	cvt.u32.u64 	%r25, %rd143;
+	cvt.u32.u64 	%r26, %rd2;
+	div.u32 	%r27, %r26, %r25;
+	cvt.u64.u32 	%rd667, %r27;
+$L__BB0_6:
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	or.b32 	%r13, %r9, 4;
+	cvt.s64.s32 	%rd4, %r12;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd151, %rd3, %rd143;
+	and.b64 	%rd152, %rd151, -4294967296;
+	setp.ne.b64 	%p27, %rd152, 0;
+	@%p27 bra 	$L__BB0_8;
+	bra.uni 	$L__BB0_7;
+$L__BB0_8:
+	div.s64 	%rd668, %rd3, %rd143;
+	bra.uni 	$L__BB0_9;
+$L__BB0_7:
+	cvt.u32.u64 	%r28, %rd143;
+	cvt.u32.u64 	%r29, %rd3;
+	div.u32 	%r30, %r29, %r28;
+	cvt.u64.u32 	%rd668, %r30;
+$L__BB0_9:
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	or.b32 	%r14, %r9, 5;
+	cvt.s64.s32 	%rd5, %r13;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd153, %rd4, %rd143;
+	and.b64 	%rd154, %rd153, -4294967296;
+	setp.ne.b64 	%p28, %rd154, 0;
+	@%p28 bra 	$L__BB0_11;
+	bra.uni 	$L__BB0_10;
+$L__BB0_11:
+	div.s64 	%rd669, %rd4, %rd143;
+	bra.uni 	$L__BB0_12;
+$L__BB0_10:
+	cvt.u32.u64 	%r31, %rd143;
+	cvt.u32.u64 	%r32, %rd4;
+	div.u32 	%r33, %r32, %r31;
+	cvt.u64.u32 	%rd669, %r33;
+$L__BB0_12:
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	or.b32 	%r15, %r9, 6;
+	cvt.s64.s32 	%rd6, %r14;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd155, %rd5, %rd143;
+	and.b64 	%rd156, %rd155, -4294967296;
+	setp.ne.b64 	%p29, %rd156, 0;
+	@%p29 bra 	$L__BB0_14;
+	bra.uni 	$L__BB0_13;
+$L__BB0_14:
+	div.s64 	%rd670, %rd5, %rd143;
+	bra.uni 	$L__BB0_15;
+$L__BB0_13:
+	cvt.u32.u64 	%r34, %rd143;
+	cvt.u32.u64 	%r35, %rd5;
+	div.u32 	%r36, %r35, %r34;
+	cvt.u64.u32 	%rd670, %r36;
+$L__BB0_15:
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	or.b32 	%r16, %r9, 7;
+	cvt.s64.s32 	%rd7, %r15;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd157, %rd6, %rd143;
+	and.b64 	%rd158, %rd157, -4294967296;
+	setp.ne.b64 	%p30, %rd158, 0;
+	@%p30 bra 	$L__BB0_17;
+	bra.uni 	$L__BB0_16;
+$L__BB0_17:
+	div.s64 	%rd671, %rd6, %rd143;
+	bra.uni 	$L__BB0_18;
+$L__BB0_16:
+	cvt.u32.u64 	%r37, %rd143;
+	cvt.u32.u64 	%r38, %rd6;
+	div.u32 	%r39, %r38, %r37;
+	cvt.u64.u32 	%rd671, %r39;
+$L__BB0_18:
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	cvt.s64.s32 	%rd8, %r16;
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd159, %rd7, %rd143;
+	and.b64 	%rd160, %rd159, -4294967296;
+	setp.ne.b64 	%p31, %rd160, 0;
+	@%p31 bra 	$L__BB0_20;
+	bra.uni 	$L__BB0_19;
+$L__BB0_20:
+	div.s64 	%rd672, %rd7, %rd143;
+	bra.uni 	$L__BB0_21;
+$L__BB0_19:
+	cvt.u32.u64 	%r40, %rd143;
+	cvt.u32.u64 	%r41, %rd7;
+	div.u32 	%r42, %r41, %r40;
+	cvt.u64.u32 	%rd672, %r42;
+$L__BB0_21:
+	.loc	1 0 21                          // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0:21
+	ld.param.b64 	%rd144, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_6];
+	.loc	1 24 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:21
+	or.b64 	%rd161, %rd8, %rd143;
+	and.b64 	%rd162, %rd161, -4294967296;
+	setp.ne.b64 	%p32, %rd162, 0;
+	@%p32 bra 	$L__BB0_23;
+	bra.uni 	$L__BB0_22;
+$L__BB0_23:
+	div.s64 	%rd673, %rd8, %rd143;
+	bra.uni 	$L__BB0_24;
+$L__BB0_22:
+	cvt.u32.u64 	%r43, %rd143;
+	cvt.u32.u64 	%r44, %rd8;
+	div.u32 	%r45, %r44, %r43;
+	cvt.u64.u32 	%rd673, %r45;
+$L__BB0_24:
+	.loc	1 24 28                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:28
+	or.b64 	%rd163, %rd666, %rd144;
+	and.b64 	%rd164, %rd163, -4294967296;
+	setp.ne.b64 	%p33, %rd164, 0;
+	@%p33 bra 	$L__BB0_26;
+	bra.uni 	$L__BB0_25;
+$L__BB0_26:
+	rem.s64 	%rd674, %rd666, %rd144;
+	bra.uni 	$L__BB0_27;
+$L__BB0_25:
+	cvt.u32.u64 	%r46, %rd144;
+	cvt.u32.u64 	%r47, %rd666;
+	rem.u32 	%r48, %r47, %r46;
+	cvt.u64.u32 	%rd674, %r48;
+$L__BB0_27:
+	or.b64 	%rd165, %rd667, %rd144;
+	and.b64 	%rd166, %rd165, -4294967296;
+	setp.ne.b64 	%p34, %rd166, 0;
+	@%p34 bra 	$L__BB0_29;
+	bra.uni 	$L__BB0_28;
+$L__BB0_29:
+	rem.s64 	%rd675, %rd667, %rd144;
+	bra.uni 	$L__BB0_30;
+$L__BB0_28:
+	cvt.u32.u64 	%r49, %rd144;
+	cvt.u32.u64 	%r50, %rd667;
+	rem.u32 	%r51, %r50, %r49;
+	cvt.u64.u32 	%rd675, %r51;
+$L__BB0_30:
+	or.b64 	%rd167, %rd668, %rd144;
+	and.b64 	%rd168, %rd167, -4294967296;
+	setp.ne.b64 	%p35, %rd168, 0;
+	@%p35 bra 	$L__BB0_32;
+	bra.uni 	$L__BB0_31;
+$L__BB0_32:
+	rem.s64 	%rd676, %rd668, %rd144;
+	bra.uni 	$L__BB0_33;
+$L__BB0_31:
+	cvt.u32.u64 	%r52, %rd144;
+	cvt.u32.u64 	%r53, %rd668;
+	rem.u32 	%r54, %r53, %r52;
+	cvt.u64.u32 	%rd676, %r54;
+$L__BB0_33:
+	or.b64 	%rd169, %rd669, %rd144;
+	and.b64 	%rd170, %rd169, -4294967296;
+	setp.ne.b64 	%p36, %rd170, 0;
+	@%p36 bra 	$L__BB0_35;
+	bra.uni 	$L__BB0_34;
+$L__BB0_35:
+	rem.s64 	%rd677, %rd669, %rd144;
+	bra.uni 	$L__BB0_36;
+$L__BB0_34:
+	cvt.u32.u64 	%r55, %rd144;
+	cvt.u32.u64 	%r56, %rd669;
+	rem.u32 	%r57, %r56, %r55;
+	cvt.u64.u32 	%rd677, %r57;
+$L__BB0_36:
+	or.b64 	%rd171, %rd670, %rd144;
+	and.b64 	%rd172, %rd171, -4294967296;
+	setp.ne.b64 	%p37, %rd172, 0;
+	@%p37 bra 	$L__BB0_38;
+	bra.uni 	$L__BB0_37;
+$L__BB0_38:
+	rem.s64 	%rd678, %rd670, %rd144;
+	bra.uni 	$L__BB0_39;
+$L__BB0_37:
+	cvt.u32.u64 	%r58, %rd144;
+	cvt.u32.u64 	%r59, %rd670;
+	rem.u32 	%r60, %r59, %r58;
+	cvt.u64.u32 	%rd678, %r60;
+$L__BB0_39:
+	or.b64 	%rd173, %rd671, %rd144;
+	and.b64 	%rd174, %rd173, -4294967296;
+	setp.ne.b64 	%p38, %rd174, 0;
+	@%p38 bra 	$L__BB0_41;
+	bra.uni 	$L__BB0_40;
+$L__BB0_41:
+	rem.s64 	%rd679, %rd671, %rd144;
+	bra.uni 	$L__BB0_42;
+$L__BB0_40:
+	cvt.u32.u64 	%r61, %rd144;
+	cvt.u32.u64 	%r62, %rd671;
+	rem.u32 	%r63, %r62, %r61;
+	cvt.u64.u32 	%rd679, %r63;
+$L__BB0_42:
+	or.b64 	%rd175, %rd672, %rd144;
+	and.b64 	%rd176, %rd175, -4294967296;
+	setp.ne.b64 	%p39, %rd176, 0;
+	@%p39 bra 	$L__BB0_44;
+	bra.uni 	$L__BB0_43;
+$L__BB0_44:
+	rem.s64 	%rd680, %rd672, %rd144;
+	bra.uni 	$L__BB0_45;
+$L__BB0_43:
+	cvt.u32.u64 	%r64, %rd144;
+	cvt.u32.u64 	%r65, %rd672;
+	rem.u32 	%r66, %r65, %r64;
+	cvt.u64.u32 	%rd680, %r66;
+$L__BB0_45:
+	.loc	1 0 28                          // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0:28
+	ld.param.b64 	%rd145, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_7];
+	ld.param.b64 	%rd139, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_1];
+	ld.param.b64 	%rd138, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_0];
+	ld.param.b32 	%r8, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_9];
+	rem.s64 	%rd9, %rd1, %rd143;
+	rem.s64 	%rd16, %rd8, %rd143;
+	rem.s64 	%rd15, %rd7, %rd143;
+	rem.s64 	%rd14, %rd6, %rd143;
+	rem.s64 	%rd13, %rd5, %rd143;
+	rem.s64 	%rd12, %rd4, %rd143;
+	rem.s64 	%rd11, %rd3, %rd143;
+	rem.s64 	%rd10, %rd2, %rd143;
+	.loc	1 24 28                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:24:28
+	or.b64 	%rd177, %rd673, %rd144;
+	and.b64 	%rd178, %rd177, -4294967296;
+	setp.ne.b64 	%p40, %rd178, 0;
+	@%p40 bra 	$L__BB0_47;
+	bra.uni 	$L__BB0_46;
+$L__BB0_47:
+	rem.s64 	%rd681, %rd673, %rd144;
+	bra.uni 	$L__BB0_48;
+$L__BB0_46:
+	cvt.u32.u64 	%r67, %rd144;
+	cvt.u32.u64 	%r68, %rd673;
+	rem.u32 	%r69, %r68, %r67;
+	cvt.u64.u32 	%rd681, %r69;
+$L__BB0_48:
+	.loc	1 21 21                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:21:21
+	setp.lt.s32 	%p8, %r16, %r8;
+	setp.lt.s32 	%p7, %r15, %r8;
+	setp.lt.s32 	%p6, %r14, %r8;
+	setp.lt.s32 	%p5, %r13, %r8;
+	setp.lt.s32 	%p4, %r12, %r8;
+	setp.lt.s32 	%p3, %r11, %r8;
+	setp.lt.s32 	%p2, %r10, %r8;
+	setp.lt.s32 	%p1, %r9, %r8;
+	.loc	1 25 31                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:25:31
+	shl.b64 	%rd291, %rd1, 1;
+	add.s64 	%rd180, %rd138, %rd291;
+	shl.b64 	%rd292, %rd2, 1;
+	add.s64 	%rd183, %rd138, %rd292;
+	shl.b64 	%rd293, %rd3, 1;
+	add.s64 	%rd186, %rd138, %rd293;
+	shl.b64 	%rd294, %rd4, 1;
+	add.s64 	%rd189, %rd138, %rd294;
+	shl.b64 	%rd295, %rd5, 1;
+	add.s64 	%rd192, %rd138, %rd295;
+	shl.b64 	%rd296, %rd6, 1;
+	add.s64 	%rd195, %rd138, %rd296;
+	shl.b64 	%rd297, %rd7, 1;
+	add.s64 	%rd198, %rd138, %rd297;
+	shl.b64 	%rd298, %rd8, 1;
+	add.s64 	%rd201, %rd138, %rd298;
+	.loc	1 25 36                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:25:36
+	// begin inline asm
+	mov.u64 %rd179, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd179, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs41, 0x0;
+	@%p1 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs41 }, [ %rd180 + 0 ], %rd179;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd182, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd182, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs42, 0x0;
+	@%p2 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs42 }, [ %rd183 + 0 ], %rd182;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd185, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd185, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs43, 0x0;
+	@%p3 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs43 }, [ %rd186 + 0 ], %rd185;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd188, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd188, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs44, 0x0;
+	@%p4 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs44 }, [ %rd189 + 0 ], %rd188;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd191, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd191, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs45, 0x0;
+	@%p5 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs45 }, [ %rd192 + 0 ], %rd191;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd194, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd194, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs46, 0x0;
+	@%p6 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs46 }, [ %rd195 + 0 ], %rd194;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd197, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd197, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs47, 0x0;
+	@%p7 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs47 }, [ %rd198 + 0 ], %rd197;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd200, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd200, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs48, 0x0;
+	@%p8 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs48 }, [ %rd201 + 0 ], %rd200;
+	// end inline asm
+	.loc	1 26 31                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:26:31
+	shl.b64 	%rd299, %rd674, 3;
+	add.s64 	%rd261, %rd139, %rd299;
+	shl.b64 	%rd300, %rd675, 3;
+	add.s64 	%rd265, %rd139, %rd300;
+	shl.b64 	%rd301, %rd676, 3;
+	add.s64 	%rd269, %rd139, %rd301;
+	shl.b64 	%rd302, %rd677, 3;
+	add.s64 	%rd273, %rd139, %rd302;
+	shl.b64 	%rd303, %rd678, 3;
+	add.s64 	%rd277, %rd139, %rd303;
+	shl.b64 	%rd304, %rd679, 3;
+	add.s64 	%rd281, %rd139, %rd304;
+	shl.b64 	%rd305, %rd680, 3;
+	add.s64 	%rd285, %rd139, %rd305;
+	shl.b64 	%rd306, %rd681, 3;
+	add.s64 	%rd289, %rd139, %rd306;
+	.loc	1 26 36                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:26:36
+	// begin inline asm
+	mov.u64 %rd203, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd203, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd204, 0x0;
+	@%p1 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd204 }, [ %rd261 + 0 ], %rd203;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd207, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd207, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd208, 0x0;
+	@%p2 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd208 }, [ %rd265 + 0 ], %rd207;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd211, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd211, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd212, 0x0;
+	@%p3 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd212 }, [ %rd269 + 0 ], %rd211;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd215, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd215, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd216, 0x0;
+	@%p4 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd216 }, [ %rd273 + 0 ], %rd215;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd219, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd219, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd220, 0x0;
+	@%p5 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd220 }, [ %rd277 + 0 ], %rd219;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd223, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd223, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd224, 0x0;
+	@%p6 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd224 }, [ %rd281 + 0 ], %rd223;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd227, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd227, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd228, 0x0;
+	@%p7 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd228 }, [ %rd285 + 0 ], %rd227;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd231, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd231, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd232, 0x0;
+	@%p8 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd232 }, [ %rd289 + 0 ], %rd231;
+	// end inline asm
+	.loc	1 28 18                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:28:18
+	shr.u64 	%rd307, %rd143, 63;
+	add.s64 	%rd308, %rd143, %rd307;
+	shr.s64 	%rd89, %rd308, 1;
+	.loc	1 30 35                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:35
+	sub.s64 	%rd309, %rd1, %rd89;
+	sub.s64 	%rd310, %rd2, %rd89;
+	sub.s64 	%rd311, %rd3, %rd89;
+	sub.s64 	%rd312, %rd4, %rd89;
+	sub.s64 	%rd313, %rd5, %rd89;
+	sub.s64 	%rd314, %rd6, %rd89;
+	sub.s64 	%rd315, %rd7, %rd89;
+	sub.s64 	%rd316, %rd8, %rd89;
+	.loc	1 30 30                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:30
+	shl.b64 	%rd317, %rd309, 1;
+	add.s64 	%rd236, %rd138, %rd317;
+	shl.b64 	%rd318, %rd310, 1;
+	add.s64 	%rd239, %rd138, %rd318;
+	shl.b64 	%rd319, %rd311, 1;
+	add.s64 	%rd242, %rd138, %rd319;
+	shl.b64 	%rd320, %rd312, 1;
+	add.s64 	%rd245, %rd138, %rd320;
+	shl.b64 	%rd321, %rd313, 1;
+	add.s64 	%rd248, %rd138, %rd321;
+	shl.b64 	%rd322, %rd314, 1;
+	add.s64 	%rd251, %rd138, %rd322;
+	shl.b64 	%rd323, %rd315, 1;
+	add.s64 	%rd254, %rd138, %rd323;
+	shl.b64 	%rd324, %rd316, 1;
+	add.s64 	%rd257, %rd138, %rd324;
+	.loc	1 30 53                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:53
+	// begin inline asm
+	mov.u64 %rd235, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd235, 1.0;
+	// end inline asm
+	.loc	1 29 19                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:29:19
+	setp.ge.s64 	%p73, %rd9, %rd89;
+	setp.ge.s64 	%p74, %rd10, %rd89;
+	setp.ge.s64 	%p75, %rd11, %rd89;
+	setp.ge.s64 	%p76, %rd12, %rd89;
+	setp.ge.s64 	%p77, %rd13, %rd89;
+	setp.ge.s64 	%p78, %rd14, %rd89;
+	setp.ge.s64 	%p79, %rd15, %rd89;
+	setp.ge.s64 	%p80, %rd16, %rd89;
+	.loc	1 30 60                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:60
+	and.pred 	%p16, %p8, %p80;
+	and.pred 	%p15, %p7, %p79;
+	and.pred 	%p14, %p6, %p78;
+	and.pred 	%p13, %p5, %p77;
+	and.pred 	%p12, %p4, %p76;
+	and.pred 	%p11, %p3, %p75;
+	and.pred 	%p10, %p2, %p74;
+	and.pred 	%p9, %p1, %p73;
+	mov.b16 	%rs64, 0;
+	.loc	1 30 53                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:53
+	// begin inline asm
+	mov.u16 %rs49, %rs64;
+	@%p9 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs49 }, [ %rd236 + 0 ], %rd235;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd238, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd238, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs51, %rs64;
+	@%p10 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs51 }, [ %rd239 + 0 ], %rd238;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd241, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd241, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs53, %rs64;
+	@%p11 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs53 }, [ %rd242 + 0 ], %rd241;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd244, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd244, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs55, %rs64;
+	@%p12 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs55 }, [ %rd245 + 0 ], %rd244;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd247, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd247, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs57, %rs64;
+	@%p13 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs57 }, [ %rd248 + 0 ], %rd247;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd250, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd250, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs59, %rs64;
+	@%p14 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs59 }, [ %rd251 + 0 ], %rd250;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd253, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd253, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs61, %rs64;
+	@%p15 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs61 }, [ %rd254 + 0 ], %rd253;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd256, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd256, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs63, %rs64;
+	@%p16 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs63 }, [ %rd257 + 0 ], %rd256;
+	// end inline asm
+	.loc	1 31 35                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:31:35
+	// begin inline asm
+	mov.u64 %rd259, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd259, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd260, 0x0;
+	@%p9 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd260 }, [ %rd261 + 0 ], %rd259;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd263, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd263, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd264, 0x0;
+	@%p10 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd264 }, [ %rd265 + 0 ], %rd263;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd267, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd267, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd268, 0x0;
+	@%p11 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd268 }, [ %rd269 + 0 ], %rd267;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd271, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd271, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd272, 0x0;
+	@%p12 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd272 }, [ %rd273 + 0 ], %rd271;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd275, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd275, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd276, 0x0;
+	@%p13 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd276 }, [ %rd277 + 0 ], %rd275;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd279, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd279, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd280, 0x0;
+	@%p14 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd280 }, [ %rd281 + 0 ], %rd279;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd283, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd283, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd284, 0x0;
+	@%p15 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd284 }, [ %rd285 + 0 ], %rd283;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd287, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd287, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd288, 0x0;
+	@%p16 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd288 }, [ %rd289 + 0 ], %rd287;
+	// end inline asm
+	.loc	1 35 32                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:35:32
+	shr.s64 	%rd325, %rd268, 63;
+	and.b64 	%rd326, %rd325, %rd145;
+	shr.s64 	%rd327, %rd272, 63;
+	and.b64 	%rd328, %rd327, %rd145;
+	shr.s64 	%rd329, %rd260, 63;
+	and.b64 	%rd330, %rd329, %rd145;
+	shr.s64 	%rd331, %rd264, 63;
+	and.b64 	%rd332, %rd331, %rd145;
+	shr.s64 	%rd333, %rd284, 63;
+	and.b64 	%rd334, %rd333, %rd145;
+	shr.s64 	%rd335, %rd288, 63;
+	and.b64 	%rd336, %rd335, %rd145;
+	shr.s64 	%rd337, %rd276, 63;
+	and.b64 	%rd338, %rd337, %rd145;
+	shr.s64 	%rd339, %rd280, 63;
+	and.b64 	%rd340, %rd339, %rd145;
+	add.s64 	%rd111, %rd340, %rd280;
+	add.s64 	%rd110, %rd338, %rd276;
+	add.s64 	%rd113, %rd336, %rd288;
+	add.s64 	%rd112, %rd334, %rd284;
+	add.s64 	%rd107, %rd332, %rd264;
+	add.s64 	%rd106, %rd330, %rd260;
+	add.s64 	%rd109, %rd328, %rd272;
+	add.s64 	%rd108, %rd326, %rd268;
+	.loc	1 36 28                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:28
+	setp.lt.s64 	%p81, %rd108, 0;
+	setp.lt.s64 	%p82, %rd109, 0;
+	setp.lt.s64 	%p83, %rd106, 0;
+	setp.lt.s64 	%p84, %rd107, 0;
+	setp.lt.s64 	%p85, %rd112, 0;
+	setp.lt.s64 	%p86, %rd113, 0;
+	setp.lt.s64 	%p87, %rd110, 0;
+	setp.lt.s64 	%p88, %rd111, 0;
+	.loc	1 36 98                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:98
+	setp.ge.s64 	%p89, %rd108, %rd145;
+	setp.ge.s64 	%p90, %rd109, %rd145;
+	setp.ge.s64 	%p91, %rd106, %rd145;
+	setp.ge.s64 	%p92, %rd107, %rd145;
+	setp.ge.s64 	%p93, %rd112, %rd145;
+	setp.ge.s64 	%p94, %rd113, %rd145;
+	setp.ge.s64 	%p95, %rd110, %rd145;
+	setp.ge.s64 	%p96, %rd111, %rd145;
+	.loc	1 36 64                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:64
+	or.pred 	%p97, %p88, %p96;
+	or.pred 	%p98, %p87, %p95;
+	or.pred 	%p99, %p86, %p94;
+	or.pred 	%p100, %p85, %p93;
+	or.pred 	%p101, %p84, %p92;
+	or.pred 	%p102, %p83, %p91;
+	or.pred 	%p103, %p82, %p90;
+	or.pred 	%p104, %p81, %p89;
+	.loc	1 36 106                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:106
+	and.pred 	%p105, %p11, %p104;
+	selp.b16 	%rs65, 1, 0, %p105;
+	shl.b16 	%rs66, %rs65, 2;
+	and.pred 	%p106, %p12, %p103;
+	selp.b16 	%rs67, -1, 0, %p106;
+	shl.b16 	%rs68, %rs67, 3;
+	or.b16 	%rs69, %rs68, %rs66;
+	and.pred 	%p107, %p9, %p102;
+	selp.b16 	%rs70, 1, 0, %p107;
+	and.pred 	%p108, %p10, %p101;
+	selp.b16 	%rs71, -1, 0, %p108;
+	shl.b16 	%rs72, %rs71, 1;
+	or.b16 	%rs73, %rs70, %rs72;
+	and.b16 	%rs74, %rs73, 3;
+	or.b16 	%rs75, %rs74, %rs69;
+	and.b16 	%rs76, %rs75, 15;
+	and.pred 	%p109, %p15, %p100;
+	selp.b16 	%rs77, 1, 0, %p109;
+	shl.b16 	%rs78, %rs77, 2;
+	and.pred 	%p110, %p16, %p99;
+	selp.b16 	%rs79, -1, 0, %p110;
+	shl.b16 	%rs80, %rs79, 3;
+	or.b16 	%rs81, %rs80, %rs78;
+	and.pred 	%p111, %p13, %p98;
+	selp.b16 	%rs82, 1, 0, %p111;
+	and.pred 	%p112, %p14, %p97;
+	selp.b16 	%rs83, -1, 0, %p112;
+	shl.b16 	%rs84, %rs83, 1;
+	or.b16 	%rs85, %rs82, %rs84;
+	and.b16 	%rs86, %rs85, 3;
+	or.b16 	%rs87, %rs86, %rs81;
+	shl.b16 	%rs88, %rs87, 4;
+	or.b16 	%rs89, %rs76, %rs88;
+	.loc	1 36 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:123
+	and.b16 	%rs90, %rs89, 255;
+	setp.eq.b16 	%p113, %rs90, 0;
+	@%p113 bra 	$L__BB0_50;
+	bra.uni 	$L__BB0_49;
+$L__BB0_50:
+	.loc	1 0 123                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0:123
+	ld.param.b64 	%rd140, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_2];
+	.loc	1 36 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:123
+	bar.sync 	0;
+	.loc	1 37 36                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:36
+	sub.s64 	%rd427, %rd9, %rd89;
+	sub.s64 	%rd428, %rd10, %rd89;
+	sub.s64 	%rd429, %rd11, %rd89;
+	sub.s64 	%rd430, %rd12, %rd89;
+	sub.s64 	%rd431, %rd13, %rd89;
+	sub.s64 	%rd432, %rd14, %rd89;
+	sub.s64 	%rd433, %rd15, %rd89;
+	sub.s64 	%rd434, %rd16, %rd89;
+	.loc	1 37 58                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:58
+	mul.lo.s64 	%rd435, %rd106, %rd143;
+	mul.lo.s64 	%rd436, %rd107, %rd143;
+	mul.lo.s64 	%rd437, %rd108, %rd143;
+	mul.lo.s64 	%rd438, %rd109, %rd143;
+	mul.lo.s64 	%rd439, %rd110, %rd143;
+	mul.lo.s64 	%rd440, %rd111, %rd143;
+	mul.lo.s64 	%rd441, %rd112, %rd143;
+	mul.lo.s64 	%rd442, %rd113, %rd143;
+	.loc	1 37 31                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:31
+	shl.b64 	%rd443, %rd427, 1;
+	add.s64 	%rd444, %rd140, %rd443;
+	shl.b64 	%rd445, %rd435, 1;
+	add.s64 	%rd348, %rd444, %rd445;
+	shl.b64 	%rd446, %rd428, 1;
+	add.s64 	%rd447, %rd140, %rd446;
+	shl.b64 	%rd448, %rd436, 1;
+	add.s64 	%rd351, %rd447, %rd448;
+	shl.b64 	%rd449, %rd429, 1;
+	add.s64 	%rd450, %rd140, %rd449;
+	shl.b64 	%rd451, %rd437, 1;
+	add.s64 	%rd354, %rd450, %rd451;
+	shl.b64 	%rd452, %rd430, 1;
+	add.s64 	%rd453, %rd140, %rd452;
+	shl.b64 	%rd454, %rd438, 1;
+	add.s64 	%rd357, %rd453, %rd454;
+	shl.b64 	%rd455, %rd431, 1;
+	add.s64 	%rd456, %rd140, %rd455;
+	shl.b64 	%rd457, %rd439, 1;
+	add.s64 	%rd360, %rd456, %rd457;
+	shl.b64 	%rd458, %rd432, 1;
+	add.s64 	%rd459, %rd140, %rd458;
+	shl.b64 	%rd460, %rd440, 1;
+	add.s64 	%rd363, %rd459, %rd460;
+	shl.b64 	%rd461, %rd433, 1;
+	add.s64 	%rd462, %rd140, %rd461;
+	shl.b64 	%rd463, %rd441, 1;
+	add.s64 	%rd366, %rd462, %rd463;
+	shl.b64 	%rd464, %rd434, 1;
+	add.s64 	%rd465, %rd140, %rd464;
+	shl.b64 	%rd466, %rd442, 1;
+	add.s64 	%rd369, %rd465, %rd466;
+	.loc	1 37 65                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:65
+	// begin inline asm
+	mov.u64 %rd347, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd347, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs91, %rs64;
+	@%p9 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs91 }, [ %rd348 + 0 ], %rd347;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd350, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd350, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs93, %rs64;
+	@%p10 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs93 }, [ %rd351 + 0 ], %rd350;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd353, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd353, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs95, %rs64;
+	@%p11 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs95 }, [ %rd354 + 0 ], %rd353;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd356, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd356, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs97, %rs64;
+	@%p12 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs97 }, [ %rd357 + 0 ], %rd356;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd359, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd359, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs99, %rs64;
+	@%p13 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs99 }, [ %rd360 + 0 ], %rd359;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd362, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd362, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs101, %rs64;
+	@%p14 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs101 }, [ %rd363 + 0 ], %rd362;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd365, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd365, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs103, %rs64;
+	@%p15 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs103 }, [ %rd366 + 0 ], %rd365;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd368, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd368, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs105, %rs64;
+	@%p16 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs105 }, [ %rd369 + 0 ], %rd368;
+	// end inline asm
+	.loc	1 44 19                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:44:19
+	setp.lt.s64 	%p146, %rd9, %rd89;
+	setp.lt.s64 	%p147, %rd10, %rd89;
+	setp.lt.s64 	%p148, %rd11, %rd89;
+	setp.lt.s64 	%p149, %rd12, %rd89;
+	setp.lt.s64 	%p150, %rd13, %rd89;
+	setp.lt.s64 	%p151, %rd14, %rd89;
+	setp.lt.s64 	%p152, %rd15, %rd89;
+	setp.lt.s64 	%p153, %rd16, %rd89;
+	.loc	1 45 37                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:37
+	add.s64 	%rd467, %rd143, %rd1;
+	add.s64 	%rd468, %rd143, %rd2;
+	add.s64 	%rd469, %rd143, %rd3;
+	add.s64 	%rd470, %rd143, %rd4;
+	add.s64 	%rd471, %rd143, %rd5;
+	add.s64 	%rd472, %rd143, %rd6;
+	add.s64 	%rd473, %rd143, %rd7;
+	add.s64 	%rd474, %rd143, %rd8;
+	.loc	1 45 42                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:42
+	sub.s64 	%rd475, %rd467, %rd89;
+	sub.s64 	%rd476, %rd468, %rd89;
+	sub.s64 	%rd477, %rd469, %rd89;
+	sub.s64 	%rd478, %rd470, %rd89;
+	sub.s64 	%rd479, %rd471, %rd89;
+	sub.s64 	%rd480, %rd472, %rd89;
+	sub.s64 	%rd481, %rd473, %rd89;
+	sub.s64 	%rd482, %rd474, %rd89;
+	.loc	1 45 31                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:31
+	shl.b64 	%rd483, %rd475, 1;
+	add.s64 	%rd372, %rd138, %rd483;
+	shl.b64 	%rd484, %rd476, 1;
+	add.s64 	%rd375, %rd138, %rd484;
+	shl.b64 	%rd485, %rd477, 1;
+	add.s64 	%rd378, %rd138, %rd485;
+	shl.b64 	%rd486, %rd478, 1;
+	add.s64 	%rd381, %rd138, %rd486;
+	shl.b64 	%rd487, %rd479, 1;
+	add.s64 	%rd384, %rd138, %rd487;
+	shl.b64 	%rd488, %rd480, 1;
+	add.s64 	%rd387, %rd138, %rd488;
+	shl.b64 	%rd489, %rd481, 1;
+	add.s64 	%rd390, %rd138, %rd489;
+	shl.b64 	%rd490, %rd482, 1;
+	add.s64 	%rd393, %rd138, %rd490;
+	.loc	1 45 60                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:60
+	// begin inline asm
+	mov.u64 %rd371, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd371, 1.0;
+	// end inline asm
+	.loc	1 45 68                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:68
+	and.pred 	%p24, %p8, %p153;
+	and.pred 	%p23, %p7, %p152;
+	and.pred 	%p22, %p6, %p151;
+	and.pred 	%p21, %p5, %p150;
+	and.pred 	%p20, %p4, %p149;
+	and.pred 	%p19, %p3, %p148;
+	and.pred 	%p18, %p2, %p147;
+	and.pred 	%p17, %p1, %p146;
+	.loc	1 45 60                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:60
+	// begin inline asm
+	mov.u16 %rs107, %rs64;
+	@%p17 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs107 }, [ %rd372 + 0 ], %rd371;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd374, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd374, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs109, %rs64;
+	@%p18 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs109 }, [ %rd375 + 0 ], %rd374;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd377, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd377, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs111, %rs64;
+	@%p19 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs111 }, [ %rd378 + 0 ], %rd377;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd380, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd380, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs113, %rs64;
+	@%p20 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs113 }, [ %rd381 + 0 ], %rd380;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd383, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd383, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs115, %rs64;
+	@%p21 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs115 }, [ %rd384 + 0 ], %rd383;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd386, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd386, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs117, %rs64;
+	@%p22 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs117 }, [ %rd387 + 0 ], %rd386;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd389, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd389, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs119, %rs64;
+	@%p23 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs119 }, [ %rd390 + 0 ], %rd389;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd392, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd392, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs121, %rs64;
+	@%p24 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs121 }, [ %rd393 + 0 ], %rd392;
+	// end inline asm
+	.loc	1 46 36                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:46:36
+	// begin inline asm
+	mov.u64 %rd395, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd395, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd396, 0x0;
+	@%p17 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd396 }, [ %rd261 + 0 ], %rd395;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd399, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd399, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd400, 0x0;
+	@%p18 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd400 }, [ %rd265 + 0 ], %rd399;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd403, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd403, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd404, 0x0;
+	@%p19 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd404 }, [ %rd269 + 0 ], %rd403;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd407, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd407, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd408, 0x0;
+	@%p20 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd408 }, [ %rd273 + 0 ], %rd407;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd411, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd411, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd412, 0x0;
+	@%p21 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd412 }, [ %rd277 + 0 ], %rd411;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd415, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd415, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd416, 0x0;
+	@%p22 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd416 }, [ %rd281 + 0 ], %rd415;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd419, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd419, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd420, 0x0;
+	@%p23 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd420 }, [ %rd285 + 0 ], %rd419;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd423, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd423, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd424, 0x0;
+	@%p24 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd424 }, [ %rd289 + 0 ], %rd423;
+	// end inline asm
+	.loc	1 50 35                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:50:35
+	shr.s64 	%rd491, %rd404, 63;
+	and.b64 	%rd492, %rd491, %rd145;
+	shr.s64 	%rd493, %rd408, 63;
+	and.b64 	%rd494, %rd493, %rd145;
+	shr.s64 	%rd495, %rd396, 63;
+	and.b64 	%rd496, %rd495, %rd145;
+	shr.s64 	%rd497, %rd400, 63;
+	and.b64 	%rd498, %rd497, %rd145;
+	shr.s64 	%rd499, %rd420, 63;
+	and.b64 	%rd500, %rd499, %rd145;
+	shr.s64 	%rd501, %rd424, 63;
+	and.b64 	%rd502, %rd501, %rd145;
+	shr.s64 	%rd503, %rd412, 63;
+	and.b64 	%rd504, %rd503, %rd145;
+	shr.s64 	%rd505, %rd416, 63;
+	and.b64 	%rd506, %rd505, %rd145;
+	add.s64 	%rd127, %rd506, %rd416;
+	add.s64 	%rd126, %rd504, %rd412;
+	add.s64 	%rd129, %rd502, %rd424;
+	add.s64 	%rd128, %rd500, %rd420;
+	add.s64 	%rd123, %rd498, %rd400;
+	add.s64 	%rd122, %rd496, %rd396;
+	add.s64 	%rd125, %rd494, %rd408;
+	add.s64 	%rd124, %rd492, %rd404;
+	.loc	1 51 28                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:28
+	setp.lt.s64 	%p154, %rd124, 0;
+	setp.lt.s64 	%p155, %rd125, 0;
+	setp.lt.s64 	%p156, %rd122, 0;
+	setp.lt.s64 	%p157, %rd123, 0;
+	setp.lt.s64 	%p158, %rd128, 0;
+	setp.lt.s64 	%p159, %rd129, 0;
+	setp.lt.s64 	%p160, %rd126, 0;
+	setp.lt.s64 	%p161, %rd127, 0;
+	.loc	1 51 100                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:100
+	setp.ge.s64 	%p162, %rd124, %rd145;
+	setp.ge.s64 	%p163, %rd125, %rd145;
+	setp.ge.s64 	%p164, %rd122, %rd145;
+	setp.ge.s64 	%p165, %rd123, %rd145;
+	setp.ge.s64 	%p166, %rd128, %rd145;
+	setp.ge.s64 	%p167, %rd129, %rd145;
+	setp.ge.s64 	%p168, %rd126, %rd145;
+	setp.ge.s64 	%p169, %rd127, %rd145;
+	.loc	1 51 65                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:65
+	or.pred 	%p170, %p161, %p169;
+	or.pred 	%p171, %p160, %p168;
+	or.pred 	%p172, %p159, %p167;
+	or.pred 	%p173, %p158, %p166;
+	or.pred 	%p174, %p157, %p165;
+	or.pred 	%p175, %p156, %p164;
+	or.pred 	%p176, %p155, %p163;
+	or.pred 	%p177, %p154, %p162;
+	.loc	1 51 108                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:108
+	and.pred 	%p178, %p19, %p177;
+	selp.b16 	%rs123, 1, 0, %p178;
+	shl.b16 	%rs124, %rs123, 2;
+	and.pred 	%p179, %p20, %p176;
+	selp.b16 	%rs125, -1, 0, %p179;
+	shl.b16 	%rs126, %rs125, 3;
+	or.b16 	%rs127, %rs126, %rs124;
+	and.pred 	%p180, %p17, %p175;
+	selp.b16 	%rs128, 1, 0, %p180;
+	and.pred 	%p181, %p18, %p174;
+	selp.b16 	%rs129, -1, 0, %p181;
+	shl.b16 	%rs130, %rs129, 1;
+	or.b16 	%rs131, %rs128, %rs130;
+	and.b16 	%rs132, %rs131, 3;
+	or.b16 	%rs133, %rs132, %rs127;
+	and.b16 	%rs134, %rs133, 15;
+	and.pred 	%p182, %p23, %p173;
+	selp.b16 	%rs135, 1, 0, %p182;
+	shl.b16 	%rs136, %rs135, 2;
+	and.pred 	%p183, %p24, %p172;
+	selp.b16 	%rs137, -1, 0, %p183;
+	shl.b16 	%rs138, %rs137, 3;
+	or.b16 	%rs139, %rs138, %rs136;
+	and.pred 	%p184, %p21, %p171;
+	selp.b16 	%rs140, 1, 0, %p184;
+	and.pred 	%p185, %p22, %p170;
+	selp.b16 	%rs141, -1, 0, %p185;
+	shl.b16 	%rs142, %rs141, 1;
+	or.b16 	%rs143, %rs140, %rs142;
+	and.b16 	%rs144, %rs143, 3;
+	or.b16 	%rs145, %rs144, %rs139;
+	shl.b16 	%rs146, %rs145, 4;
+	or.b16 	%rs147, %rs134, %rs146;
+	.loc	1 51 126                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:126
+	and.b16 	%rs148, %rs147, 255;
+	setp.eq.b16 	%p186, %rs148, 0;
+	@%p186 bra 	$L__BB0_52;
+	bra.uni 	$L__BB0_51;
+$L__BB0_52:
+	.loc	1 0 126                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0:126
+	ld.param.b64 	%rd146, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_8];
+	.loc	1 51 126                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:126
+	bar.sync 	0;
+	.loc	1 52 37                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:37
+	sub.s64 	%rd537, %rd143, %rd89;
+	.loc	1 52 64                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:64
+	mul.lo.s64 	%rd538, %rd122, %rd143;
+	mul.lo.s64 	%rd539, %rd123, %rd143;
+	mul.lo.s64 	%rd540, %rd124, %rd143;
+	mul.lo.s64 	%rd541, %rd125, %rd143;
+	mul.lo.s64 	%rd542, %rd126, %rd143;
+	mul.lo.s64 	%rd543, %rd127, %rd143;
+	mul.lo.s64 	%rd544, %rd128, %rd143;
+	mul.lo.s64 	%rd545, %rd129, %rd143;
+	.loc	1 52 31                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:31
+	shl.b64 	%rd546, %rd537, 1;
+	add.s64 	%rd547, %rd140, %rd546;
+	shl.b64 	%rd548, %rd9, 1;
+	add.s64 	%rd549, %rd547, %rd548;
+	shl.b64 	%rd550, %rd538, 1;
+	add.s64 	%rd514, %rd549, %rd550;
+	shl.b64 	%rd551, %rd10, 1;
+	add.s64 	%rd552, %rd547, %rd551;
+	shl.b64 	%rd553, %rd539, 1;
+	add.s64 	%rd517, %rd552, %rd553;
+	shl.b64 	%rd554, %rd11, 1;
+	add.s64 	%rd555, %rd547, %rd554;
+	shl.b64 	%rd556, %rd540, 1;
+	add.s64 	%rd520, %rd555, %rd556;
+	shl.b64 	%rd557, %rd12, 1;
+	add.s64 	%rd558, %rd547, %rd557;
+	shl.b64 	%rd559, %rd541, 1;
+	add.s64 	%rd523, %rd558, %rd559;
+	shl.b64 	%rd560, %rd13, 1;
+	add.s64 	%rd561, %rd547, %rd560;
+	shl.b64 	%rd562, %rd542, 1;
+	add.s64 	%rd526, %rd561, %rd562;
+	shl.b64 	%rd563, %rd14, 1;
+	add.s64 	%rd564, %rd547, %rd563;
+	shl.b64 	%rd565, %rd543, 1;
+	add.s64 	%rd529, %rd564, %rd565;
+	shl.b64 	%rd566, %rd15, 1;
+	add.s64 	%rd567, %rd547, %rd566;
+	shl.b64 	%rd568, %rd544, 1;
+	add.s64 	%rd532, %rd567, %rd568;
+	shl.b64 	%rd569, %rd16, 1;
+	add.s64 	%rd570, %rd547, %rd569;
+	shl.b64 	%rd571, %rd545, 1;
+	add.s64 	%rd535, %rd570, %rd571;
+	.loc	1 52 72                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:72
+	// begin inline asm
+	mov.u64 %rd513, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd513, 1.0;
+	// end inline asm
+	mov.b16 	%rs150, 0;
+	// begin inline asm
+	mov.u16 %rs149, %rs150;
+	@%p17 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs149 }, [ %rd514 + 0 ], %rd513;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd516, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd516, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs151, %rs150;
+	@%p18 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs151 }, [ %rd517 + 0 ], %rd516;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd519, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd519, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs153, %rs150;
+	@%p19 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs153 }, [ %rd520 + 0 ], %rd519;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd522, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd522, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs155, %rs150;
+	@%p20 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs155 }, [ %rd523 + 0 ], %rd522;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd525, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd525, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs157, %rs150;
+	@%p21 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs157 }, [ %rd526 + 0 ], %rd525;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd528, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd528, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs159, %rs150;
+	@%p22 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs159 }, [ %rd529 + 0 ], %rd528;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd531, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd531, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs161, %rs150;
+	@%p23 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs161 }, [ %rd532 + 0 ], %rd531;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd534, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd534, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs163, %rs150;
+	@%p24 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs163 }, [ %rd535 + 0 ], %rd534;
+	// end inline asm
+	.loc	1 61 35                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:61:35
+	shr.s64 	%rd572, %rd212, 63;
+	and.b64 	%rd573, %rd572, %rd146;
+	shr.s64 	%rd574, %rd216, 63;
+	and.b64 	%rd575, %rd574, %rd146;
+	shr.s64 	%rd576, %rd204, 63;
+	and.b64 	%rd577, %rd576, %rd146;
+	shr.s64 	%rd578, %rd208, 63;
+	and.b64 	%rd579, %rd578, %rd146;
+	shr.s64 	%rd580, %rd228, 63;
+	and.b64 	%rd581, %rd580, %rd146;
+	shr.s64 	%rd582, %rd232, 63;
+	and.b64 	%rd583, %rd582, %rd146;
+	shr.s64 	%rd584, %rd220, 63;
+	and.b64 	%rd585, %rd584, %rd146;
+	shr.s64 	%rd586, %rd224, 63;
+	and.b64 	%rd587, %rd586, %rd146;
+	add.s64 	%rd135, %rd587, %rd224;
+	add.s64 	%rd134, %rd585, %rd220;
+	add.s64 	%rd137, %rd583, %rd232;
+	add.s64 	%rd136, %rd581, %rd228;
+	add.s64 	%rd131, %rd579, %rd208;
+	add.s64 	%rd130, %rd577, %rd204;
+	add.s64 	%rd133, %rd575, %rd216;
+	add.s64 	%rd132, %rd573, %rd212;
+	.loc	1 62 28                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:28
+	setp.lt.s64 	%p203, %rd132, 0;
+	setp.lt.s64 	%p204, %rd133, 0;
+	setp.lt.s64 	%p205, %rd130, 0;
+	setp.lt.s64 	%p206, %rd131, 0;
+	setp.lt.s64 	%p207, %rd136, 0;
+	setp.lt.s64 	%p208, %rd137, 0;
+	setp.lt.s64 	%p209, %rd134, 0;
+	setp.lt.s64 	%p210, %rd135, 0;
+	.loc	1 62 46                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:46
+	setp.ge.s64 	%p211, %rd132, %rd146;
+	setp.ge.s64 	%p212, %rd133, %rd146;
+	setp.ge.s64 	%p213, %rd130, %rd146;
+	setp.ge.s64 	%p214, %rd131, %rd146;
+	setp.ge.s64 	%p215, %rd136, %rd146;
+	setp.ge.s64 	%p216, %rd137, %rd146;
+	setp.ge.s64 	%p217, %rd134, %rd146;
+	setp.ge.s64 	%p218, %rd135, %rd146;
+	.loc	1 62 38                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:38
+	or.pred 	%p219, %p210, %p218;
+	or.pred 	%p220, %p209, %p217;
+	or.pred 	%p221, %p208, %p216;
+	or.pred 	%p222, %p207, %p215;
+	or.pred 	%p223, %p206, %p214;
+	or.pred 	%p224, %p205, %p213;
+	or.pred 	%p225, %p204, %p212;
+	or.pred 	%p226, %p203, %p211;
+	.loc	1 62 54                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:54
+	and.pred 	%p227, %p3, %p226;
+	selp.b16 	%rs165, 1, 0, %p227;
+	shl.b16 	%rs166, %rs165, 2;
+	and.pred 	%p228, %p4, %p225;
+	selp.b16 	%rs167, -1, 0, %p228;
+	shl.b16 	%rs168, %rs167, 3;
+	or.b16 	%rs169, %rs168, %rs166;
+	and.pred 	%p229, %p1, %p224;
+	selp.b16 	%rs170, 1, 0, %p229;
+	and.pred 	%p230, %p2, %p223;
+	selp.b16 	%rs171, -1, 0, %p230;
+	shl.b16 	%rs172, %rs171, 1;
+	or.b16 	%rs173, %rs170, %rs172;
+	and.b16 	%rs174, %rs173, 3;
+	or.b16 	%rs175, %rs174, %rs169;
+	and.b16 	%rs176, %rs175, 15;
+	and.pred 	%p231, %p7, %p222;
+	selp.b16 	%rs177, 1, 0, %p231;
+	shl.b16 	%rs178, %rs177, 2;
+	and.pred 	%p232, %p8, %p221;
+	selp.b16 	%rs179, -1, 0, %p232;
+	shl.b16 	%rs180, %rs179, 3;
+	or.b16 	%rs181, %rs180, %rs178;
+	and.pred 	%p233, %p5, %p220;
+	selp.b16 	%rs182, 1, 0, %p233;
+	and.pred 	%p234, %p6, %p219;
+	selp.b16 	%rs183, -1, 0, %p234;
+	shl.b16 	%rs184, %rs183, 1;
+	or.b16 	%rs185, %rs182, %rs184;
+	and.b16 	%rs186, %rs185, 3;
+	or.b16 	%rs187, %rs186, %rs181;
+	shl.b16 	%rs188, %rs187, 4;
+	or.b16 	%rs189, %rs176, %rs188;
+	.loc	1 62 64                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:64
+	and.b16 	%rs190, %rs189, 255;
+	setp.eq.b16 	%p235, %rs190, 0;
+	@%p235 bra 	$L__BB0_54;
+	bra.uni 	$L__BB0_53;
+$L__BB0_54:
+	.loc	1 0 64                          // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0:64
+	ld.param.b64 	%rd142, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_4];
+	ld.param.b64 	%rd141, [triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0_param_3];
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r70, %rs63;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r71, %rs105;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r72, %r70;
+	fma.rn.f32 	%r73, %r72, %r71, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r74, %r73, 0f00000000, %p80;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r75, %rs121;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r76, %rs163;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r77, %r75, %r76;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r78, %r77, 0f00000000, %p153;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r79, %r74, %r78;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r80, %rs61;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r81, %rs103;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r82, %r80;
+	fma.rn.f32 	%r83, %r82, %r81, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r84, %r83, 0f00000000, %p79;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r85, %rs119;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r86, %rs161;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r87, %r85, %r86;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r88, %r87, 0f00000000, %p152;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r89, %r84, %r88;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r90, %rs59;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r91, %rs101;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r92, %r90;
+	fma.rn.f32 	%r93, %r92, %r91, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r94, %r93, 0f00000000, %p78;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r95, %rs117;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r96, %rs159;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r97, %r95, %r96;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r98, %r97, 0f00000000, %p151;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r99, %r94, %r98;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r100, %rs57;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r101, %rs99;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r102, %r100;
+	fma.rn.f32 	%r103, %r102, %r101, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r104, %r103, 0f00000000, %p77;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r105, %rs115;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r106, %rs157;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r107, %r105, %r106;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r108, %r107, 0f00000000, %p150;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r109, %r104, %r108;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r110, %rs55;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r111, %rs97;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r112, %r110;
+	fma.rn.f32 	%r113, %r112, %r111, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r114, %r113, 0f00000000, %p76;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r115, %rs113;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r116, %rs155;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r117, %r115, %r116;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r118, %r117, 0f00000000, %p149;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r119, %r114, %r118;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r120, %rs53;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r121, %rs95;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r122, %r120;
+	fma.rn.f32 	%r123, %r122, %r121, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r124, %r123, 0f00000000, %p75;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r125, %rs111;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r126, %rs153;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r127, %r125, %r126;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r128, %r127, 0f00000000, %p148;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r129, %r124, %r128;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r130, %rs51;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r131, %rs93;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r132, %r130;
+	fma.rn.f32 	%r133, %r132, %r131, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r134, %r133, 0f00000000, %p74;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r135, %rs109;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r136, %rs151;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r137, %r135, %r136;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r138, %r137, 0f00000000, %p147;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r139, %r134, %r138;
+	.loc	1 30 111                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:30:111
+	cvt.f32.bf16 	%r140, %rs49;
+	.loc	1 37 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:37:123
+	cvt.f32.bf16 	%r141, %rs91;
+	.loc	1 39 13                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:39:13
+	neg.f32 	%r142, %r140;
+	fma.rn.f32 	%r143, %r142, %r141, 0f00000000;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r144, %r143, 0f00000000, %p73;
+	.loc	1 45 119                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:45:119
+	cvt.f32.bf16 	%r145, %rs107;
+	.loc	1 52 131                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:52:131
+	cvt.f32.bf16 	%r146, %rs149;
+	.loc	1 53 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:53:20
+	mul.f32 	%r147, %r145, %r146;
+	.loc	1 0 0                           // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:0
+	selp.f32 	%r148, %r147, 0f00000000, %p146;
+	.loc	1 57 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:57:20
+	add.f32 	%r149, %r144, %r148;
+	.loc	1 25 76                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:25:76
+	cvt.f32.bf16 	%r150, %rs48;
+	cvt.f32.bf16 	%r151, %rs47;
+	cvt.f32.bf16 	%r152, %rs46;
+	cvt.f32.bf16 	%r153, %rs45;
+	cvt.f32.bf16 	%r154, %rs44;
+	cvt.f32.bf16 	%r155, %rs43;
+	cvt.f32.bf16 	%r156, %rs42;
+	cvt.f32.bf16 	%r157, %rs41;
+	.loc	1 62 64                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:64
+	bar.sync 	0;
+	.loc	1 63 40                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:63:40
+	mul.lo.s64 	%rd626, %rd130, %rd143;
+	mul.lo.s64 	%rd627, %rd131, %rd143;
+	mul.lo.s64 	%rd628, %rd132, %rd143;
+	mul.lo.s64 	%rd629, %rd133, %rd143;
+	mul.lo.s64 	%rd630, %rd134, %rd143;
+	mul.lo.s64 	%rd631, %rd135, %rd143;
+	mul.lo.s64 	%rd632, %rd136, %rd143;
+	mul.lo.s64 	%rd633, %rd137, %rd143;
+	.loc	1 63 31                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:63:31
+	add.s64 	%rd635, %rd141, %rd548;
+	shl.b64 	%rd636, %rd626, 1;
+	add.s64 	%rd595, %rd635, %rd636;
+	add.s64 	%rd638, %rd141, %rd551;
+	shl.b64 	%rd639, %rd627, 1;
+	add.s64 	%rd598, %rd638, %rd639;
+	add.s64 	%rd641, %rd141, %rd554;
+	shl.b64 	%rd642, %rd628, 1;
+	add.s64 	%rd601, %rd641, %rd642;
+	add.s64 	%rd644, %rd141, %rd557;
+	shl.b64 	%rd645, %rd629, 1;
+	add.s64 	%rd604, %rd644, %rd645;
+	add.s64 	%rd647, %rd141, %rd560;
+	shl.b64 	%rd648, %rd630, 1;
+	add.s64 	%rd607, %rd647, %rd648;
+	add.s64 	%rd650, %rd141, %rd563;
+	shl.b64 	%rd651, %rd631, 1;
+	add.s64 	%rd610, %rd650, %rd651;
+	add.s64 	%rd653, %rd141, %rd566;
+	shl.b64 	%rd654, %rd632, 1;
+	add.s64 	%rd613, %rd653, %rd654;
+	add.s64 	%rd656, %rd141, %rd569;
+	shl.b64 	%rd657, %rd633, 1;
+	add.s64 	%rd616, %rd656, %rd657;
+	.loc	1 63 48                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:63:48
+	// begin inline asm
+	mov.u64 %rd596, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd596, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs191, 0x0;
+	@%p1 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs191 }, [ %rd595 + 0 ], %rd596;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd599, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd599, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs192, 0x0;
+	@%p2 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs192 }, [ %rd598 + 0 ], %rd599;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd602, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd602, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs193, 0x0;
+	@%p3 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs193 }, [ %rd601 + 0 ], %rd602;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd605, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd605, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs194, 0x0;
+	@%p4 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs194 }, [ %rd604 + 0 ], %rd605;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd608, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd608, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs195, 0x0;
+	@%p5 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs195 }, [ %rd607 + 0 ], %rd608;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd611, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd611, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs196, 0x0;
+	@%p6 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs196 }, [ %rd610 + 0 ], %rd611;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd614, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd614, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs197, 0x0;
+	@%p7 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs197 }, [ %rd613 + 0 ], %rd614;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd617, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd617, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u16 %rs198, 0x0;
+	@%p8 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs198 }, [ %rd616 + 0 ], %rd617;
+	// end inline asm
+	.loc	1 63 88                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:63:88
+	cvt.f32.bf16 	%r158, %rs191;
+	cvt.f32.bf16 	%r159, %rs192;
+	cvt.f32.bf16 	%r160, %rs193;
+	cvt.f32.bf16 	%r161, %rs194;
+	cvt.f32.bf16 	%r162, %rs195;
+	cvt.f32.bf16 	%r163, %rs196;
+	cvt.f32.bf16 	%r164, %rs197;
+	cvt.f32.bf16 	%r165, %rs198;
+	.loc	1 65 20                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:65:20
+	fma.rn.f32 	%r166, %r157, %r158, %r149;
+	fma.rn.f32 	%r167, %r156, %r159, %r139;
+	fma.rn.f32 	%r168, %r155, %r160, %r129;
+	fma.rn.f32 	%r169, %r154, %r161, %r119;
+	fma.rn.f32 	%r170, %r153, %r162, %r109;
+	fma.rn.f32 	%r171, %r152, %r163, %r99;
+	fma.rn.f32 	%r172, %r151, %r164, %r89;
+	fma.rn.f32 	%r173, %r150, %r165, %r79;
+	.loc	1 66 25                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:66:25
+	add.s64 	%rd618, %rd142, %rd291;
+	add.s64 	%rd619, %rd142, %rd292;
+	add.s64 	%rd620, %rd142, %rd293;
+	add.s64 	%rd621, %rd142, %rd294;
+	add.s64 	%rd622, %rd142, %rd295;
+	add.s64 	%rd623, %rd142, %rd296;
+	add.s64 	%rd624, %rd142, %rd297;
+	add.s64 	%rd625, %rd142, %rd298;
+	.loc	1 66 37                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:66:37
+	cvt.rn.bf16.f32 	%rs199, %r166;
+	cvt.rn.bf16.f32 	%rs200, %r167;
+	cvt.rn.bf16.f32 	%rs201, %r168;
+	cvt.rn.bf16.f32 	%rs202, %r169;
+	cvt.rn.bf16.f32 	%rs203, %r170;
+	cvt.rn.bf16.f32 	%rs204, %r171;
+	cvt.rn.bf16.f32 	%rs205, %r172;
+	cvt.rn.bf16.f32 	%rs206, %r173;
+	// begin inline asm
+	@%p1 st.global.b16 [ %rd618 + 0 ], { %rs199 };
+	// end inline asm
+	// begin inline asm
+	@%p2 st.global.b16 [ %rd619 + 0 ], { %rs200 };
+	// end inline asm
+	// begin inline asm
+	@%p3 st.global.b16 [ %rd620 + 0 ], { %rs201 };
+	// end inline asm
+	// begin inline asm
+	@%p4 st.global.b16 [ %rd621 + 0 ], { %rs202 };
+	// end inline asm
+	// begin inline asm
+	@%p5 st.global.b16 [ %rd622 + 0 ], { %rs203 };
+	// end inline asm
+	// begin inline asm
+	@%p6 st.global.b16 [ %rd623 + 0 ], { %rs204 };
+	// end inline asm
+	// begin inline asm
+	@%p7 st.global.b16 [ %rd624 + 0 ], { %rs205 };
+	// end inline asm
+	// begin inline asm
+	@%p8 st.global.b16 [ %rd625 + 0 ], { %rs206 };
+	// end inline asm
+	.loc	1 66 4                          // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:66:4
+	ret;
+$L__BB0_49:
+	.loc	1 36 123                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:36:123
+	{ // callseq 0, 0
+	.param .b64 	param0;
+	.param .b64 	param1;
+	.param .b32 	param2;
+	.param .b64 	param3;
+	.param .b64 	param4;
+	mov.b64 	%rd341, assertFunc_0;
+	cvta.global.u64 	%rd342, %rd341;
+	st.param.b64 	[param3], %rd342;
+	mov.b64 	%rd343, assertFile_0;
+	cvta.global.u64 	%rd344, %rd343;
+	st.param.b64 	[param1], %rd344;
+	mov.b64 	%rd345, assertMessage_0;
+	cvta.global.u64 	%rd346, %rd345;
+	st.param.b64 	[param0], %rd346;
+	st.param.b64 	[param4], 1;
+	st.param.b32 	[param2], 36;
+	call.uni __assertfail, (param0, param1, param2, param3, param4);
+	} // callseq 0
+	trap;
+$L__BB0_51:
+	.loc	1 51 126                        // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:51:126
+	{ // callseq 1, 0
+	.param .b64 	param0;
+	.param .b64 	param1;
+	.param .b32 	param2;
+	.param .b64 	param3;
+	.param .b64 	param4;
+	mov.b64 	%rd507, assertFunc_1;
+	cvta.global.u64 	%rd508, %rd507;
+	st.param.b64 	[param3], %rd508;
+	mov.b64 	%rd509, assertFile_1;
+	cvta.global.u64 	%rd510, %rd509;
+	st.param.b64 	[param1], %rd510;
+	mov.b64 	%rd511, assertMessage_1;
+	cvta.global.u64 	%rd512, %rd511;
+	st.param.b64 	[param0], %rd512;
+	st.param.b64 	[param4], 1;
+	st.param.b32 	[param2], 51;
+	call.uni __assertfail, (param0, param1, param2, param3, param4);
+	} // callseq 1
+	trap;
+$L__BB0_53:
+	.loc	1 62 64                         // cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py:62:64
+	{ // callseq 2, 0
+	.param .b64 	param0;
+	.param .b64 	param1;
+	.param .b32 	param2;
+	.param .b64 	param3;
+	.param .b64 	param4;
+	mov.b64 	%rd588, assertFunc_2;
+	cvta.global.u64 	%rd589, %rd588;
+	st.param.b64 	[param3], %rd589;
+	mov.b64 	%rd590, assertFile_2;
+	cvta.global.u64 	%rd591, %rd590;
+	st.param.b64 	[param1], %rd591;
+	mov.b64 	%rd592, assertMessage_2;
+	cvta.global.u64 	%rd593, %rd592;
+	st.param.b64 	[param0], %rd593;
+	st.param.b64 	[param4], 1;
+	st.param.b32 	[param2], 62;
+	call.uni __assertfail, (param0, param1, param2, param3, param4);
+	} // callseq 2
+	trap;
+$L__tmp1:
+$L__func_end0:
+                                        // -- End function
+}
+	.file	1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py"
+	.section	.debug_abbrev
+	{
+.b8 1                                   // Abbreviation Code
+.b8 17                                  // DW_TAG_compile_unit
+.b8 0                                   // DW_CHILDREN_no
+.b8 37                                  // DW_AT_producer
+.b8 8                                   // DW_FORM_string
+.b8 19                                  // DW_AT_language
+.b8 5                                   // DW_FORM_data2
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 16                                  // DW_AT_stmt_list
+.b8 6                                   // DW_FORM_data4
+.b8 27                                  // DW_AT_comp_dir
+.b8 8                                   // DW_FORM_string
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 0                                   // EOM(3)
+	}
+	.section	.debug_info
+	{
+.b32 135                                // Length of Unit
+.b8 2                                   // DWARF version number
+.b8 0
+.b32 .debug_abbrev                      // Offset Into Abbrev. Section
+.b8 8                                   // Address Size (in bytes)
+.b8 1                                   // Abbrev [1] 0xb:0x80 DW_TAG_compile_unit
+.b8 116                                 // DW_AT_producer
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2                                   // DW_AT_language
+.b8 0
+.b8 99                                  // DW_AT_name
+.b8 109
+.b8 98
+.b8 53
+.b8 122
+.b8 108
+.b8 108
+.b8 100
+.b8 107
+.b8 114
+.b8 102
+.b8 108
+.b8 101
+.b8 121
+.b8 104
+.b8 99
+.b8 107
+.b8 104
+.b8 105
+.b8 111
+.b8 109
+.b8 104
+.b8 112
+.b8 102
+.b8 115
+.b8 118
+.b8 107
+.b8 102
+.b8 107
+.b8 108
+.b8 108
+.b8 111
+.b8 116
+.b8 98
+.b8 111
+.b8 112
+.b8 97
+.b8 108
+.b8 100
+.b8 100
+.b8 119
+.b8 107
+.b8 100
+.b8 102
+.b8 104
+.b8 108
+.b8 105
+.b8 122
+.b8 118
+.b8 101
+.b8 117
+.b8 120
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line                        // DW_AT_stmt_list
+.b8 47                                  // DW_AT_comp_dir
+.b8 119
+.b8 111
+.b8 114
+.b8 107
+.b8 115
+.b8 112
+.b8 97
+.b8 99
+.b8 101
+.b8 47
+.b8 104
+.b8 97
+.b8 110
+.b8 114
+.b8 117
+.b8 105
+.b8 47
+.b8 83
+.b8 112
+.b8 101
+.b8 99
+.b8 70
+.b8 111
+.b8 114
+.b8 103
+.b8 101
+.b8 45
+.b8 101
+.b8 120
+.b8 116
+.b8 47
+.b8 99
+.b8 97
+.b8 99
+.b8 104
+.b8 101
+.b8 47
+.b8 99
+.b8 111
+.b8 109
+.b8 112
+.b8 105
+.b8 108
+.b8 101
+.b8 100
+.b8 95
+.b8 107
+.b8 101
+.b8 114
+.b8 110
+.b8 101
+.b8 108
+.b8 115
+.b8 47
+.b8 109
+.b8 98
+.b8 0
+	}
+	.section	.debug_macinfo	{	}

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.source

@@ -0,0 +1,419 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":18:0)
+#loc109 = loc("in_ptr0"(#loc))
+#loc110 = loc("in_ptr1"(#loc))
+#loc111 = loc("in_ptr2"(#loc))
+#loc112 = loc("in_ptr3"(#loc))
+#loc113 = loc("out_ptr0"(#loc))
+#loc114 = loc("ks0"(#loc))
+#loc115 = loc("ks1"(#loc))
+#loc116 = loc("ks2"(#loc))
+#loc117 = loc("ks3"(#loc))
+#loc118 = loc("xnumel"(#loc))
+module {
+  tt.func public @triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %in_ptr2: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr2"(#loc)), %in_ptr3: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr3"(#loc)), %out_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %ks2: i64 loc("ks2"(#loc)), %ks3: i64 loc("ks3"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
+    %xoffset = tt.get_program_id x : i32 loc(#loc119)
+    %xoffset_0 = arith.constant 1024 : i32 loc(#loc120)
+    %xoffset_1 = arith.constant 1024 : i32 loc(#loc120)
+    %xoffset_2 = arith.muli %xoffset, %xoffset_1 : i32 loc(#loc120)
+    %xindex = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc121)
+    %xindex_3 = tt.splat %xoffset_2 : i32 -> tensor<1024xi32> loc(#loc122)
+    %xindex_4 = arith.addi %xindex_3, %xindex : tensor<1024xi32> loc(#loc122)
+    %xmask = tt.splat %xnumel : i32 -> tensor<1024xi32> loc(#loc123)
+    %xmask_5 = arith.cmpi slt, %xindex_4, %xmask : tensor<1024xi32> loc(#loc123)
+    %x0 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc124)
+    %x0_6 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc124)
+    %x0_7 = arith.remsi %x0, %x0_6 : tensor<1024xi64> loc(#loc124)
+    %x1 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc125)
+    %x1_8 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc125)
+    %x1_9 = arith.divsi %x1, %x1_8 : tensor<1024xi64> loc(#loc125)
+    %x1_10 = tt.splat %ks1 : i64 -> tensor<1024xi64> loc(#loc126)
+    %x1_11 = arith.remsi %x1_9, %x1_10 : tensor<1024xi64> loc(#loc126)
+    %tmp31 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc127)
+    %tmp31_12 = tt.addptr %tmp31, %xindex_4 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc127)
+    %tmp31_13 = tt.load %tmp31_12, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc128)
+    %tmp31_14 = arith.extf %tmp31_13 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc129)
+    %tmp32 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>> loc(#loc130)
+    %tmp32_15 = tt.addptr %tmp32, %x1_11 : tensor<1024x!tt.ptr<i64>>, tensor<1024xi64> loc(#loc130)
+    %tmp32_16 = tt.load %tmp32_15, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc131)
+    %tmp1 = arith.constant 2 : i32 loc(#loc132)
+    %tmp1_17 = arith.constant 2 : i64 loc(#loc132)
+    %tmp1_18 = arith.divsi %ks0, %tmp1_17 : i64 loc(#loc132)
+    %tmp2 = tt.splat %tmp1_18 : i64 -> tensor<1024xi64> loc(#loc133)
+    %tmp2_19 = arith.cmpi sge, %x0_7, %tmp2 : tensor<1024xi64> loc(#loc133)
+    %tmp3 = arith.constant 2 : i32 loc(#loc134)
+    %tmp3_20 = arith.constant 2 : i64 loc(#loc134)
+    %tmp3_21 = arith.divsi %ks0, %tmp3_20 : i64 loc(#loc134)
+    %tmp3_22 = arith.constant -1 : i32 loc(#loc135)
+    %tmp3_23 = arith.constant -1 : i64 loc(#loc135)
+    %tmp3_24 = arith.muli %tmp3_23, %tmp3_21 : i64 loc(#loc135)
+    %tmp3_25 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc136)
+    %tmp3_26 = tt.splat %tmp3_24 : i64 -> tensor<1024xi64> loc(#loc136)
+    %tmp3_27 = arith.addi %tmp3_25, %tmp3_26 : tensor<1024xi64> loc(#loc136)
+    %tmp3_28 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc137)
+    %tmp3_29 = tt.addptr %tmp3_28, %tmp3_27 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc137)
+    %tmp3_30 = arith.andi %tmp2_19, %xmask_5 : tensor<1024xi1> loc(#loc138)
+    %tmp3_31 = arith.constant 0.000000e+00 : f32 loc(#loc139)
+    %tmp3_32 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc139)
+    %tmp3_33 = arith.truncf %tmp3_32 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc139)
+    %tmp3_34 = tt.load %tmp3_29, %tmp3_30, %tmp3_33 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc139)
+    %tmp3_35 = arith.extf %tmp3_34 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc140)
+    %tmp4 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>> loc(#loc141)
+    %tmp4_36 = tt.addptr %tmp4, %x1_11 : tensor<1024x!tt.ptr<i64>>, tensor<1024xi64> loc(#loc141)
+    %tmp4_37 = arith.andi %tmp2_19, %xmask_5 : tensor<1024xi1> loc(#loc142)
+    %tmp4_38 = arith.constant 0.000000e+00 : f32 loc(#loc143)
+    %tmp4_39 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc143)
+    %tmp4_40 = arith.fptosi %tmp4_39 : tensor<1024xf32> to tensor<1024xi64> loc(#loc143)
+    %tmp4_41 = tt.load %tmp4_36, %tmp4_37, %tmp4_40 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc143)
+    %tmp5 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc144)
+    %tmp6 = arith.addi %tmp4_41, %tmp5 : tensor<1024xi64> loc(#loc145)
+    %tmp7 = arith.constant 0 : i32 loc(#loc146)
+    %tmp7_42 = arith.extsi %tmp7 : i32 to i64 loc(#loc146)
+    %tmp7_43 = tt.splat %tmp7_42 : i64 -> tensor<1024xi64> loc(#loc146)
+    %tmp7_44 = arith.cmpi slt, %tmp4_41, %tmp7_43 : tensor<1024xi64> loc(#loc146)
+    %tmp8 = arith.select %tmp7_44, %tmp6, %tmp4_41 : tensor<1024xi1>, tensor<1024xi64> loc(#loc147)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc30)
+    %0 = arith.extsi %c0_i32 : i32 to i64 loc(#loc30)
+    %1 = tt.splat %0 : i64 -> tensor<1024xi64> loc(#loc30)
+    %2 = arith.cmpi sle, %1, %tmp8 : tensor<1024xi64> loc(#loc30)
+    %3 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc31)
+    %4 = arith.cmpi slt, %tmp8, %3 : tensor<1024xi64> loc(#loc31)
+    %5 = arith.andi %2, %4 : tensor<1024xi1> loc(#loc32)
+    %6 = arith.andi %tmp2_19, %xmask_5 : tensor<1024xi1> loc(#loc33)
+    %true = arith.constant true loc(#loc34)
+    %cst = arith.constant dense<true> : tensor<1024xi1> loc(#loc34)
+    %7 = arith.xori %6, %cst : tensor<1024xi1> loc(#loc34)
+    %8 = arith.ori %5, %7 : tensor<1024xi1> loc(#loc35)
+    tt.assert %8, "index out of bounds: 0 <= tl.broadcast_to(tmp8, [XBLOCK]) < ks2" : tensor<1024xi1> loc(#loc36)
+    %tmp10 = arith.constant 2 : i32 loc(#loc148)
+    %tmp10_45 = arith.constant 2 : i64 loc(#loc148)
+    %tmp10_46 = arith.divsi %ks0, %tmp10_45 : i64 loc(#loc148)
+    %tmp10_47 = arith.constant -1 : i32 loc(#loc149)
+    %tmp10_48 = arith.constant -1 : i64 loc(#loc149)
+    %tmp10_49 = arith.muli %tmp10_48, %tmp10_46 : i64 loc(#loc149)
+    %tmp10_50 = tt.splat %tmp10_49 : i64 -> tensor<1024xi64> loc(#loc150)
+    %tmp10_51 = arith.addi %x0_7, %tmp10_50 : tensor<1024xi64> loc(#loc150)
+    %tmp10_52 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc151)
+    %tmp10_53 = arith.muli %tmp10_52, %tmp8 : tensor<1024xi64> loc(#loc151)
+    %tmp10_54 = arith.addi %tmp10_51, %tmp10_53 : tensor<1024xi64> loc(#loc152)
+    %tmp10_55 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc153)
+    %tmp10_56 = tt.addptr %tmp10_55, %tmp10_54 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc153)
+    %tmp10_57 = arith.andi %tmp2_19, %xmask_5 : tensor<1024xi1> loc(#loc154)
+    %tmp10_58 = arith.constant 0.000000e+00 : f32 loc(#loc155)
+    %tmp10_59 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc155)
+    %tmp10_60 = arith.truncf %tmp10_59 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc155)
+    %tmp10_61 = tt.load %tmp10_56, %tmp10_57, %tmp10_60 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc155)
+    %tmp10_62 = arith.extf %tmp10_61 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc156)
+    %tmp11 = arith.mulf %tmp3_35, %tmp10_62 : tensor<1024xf32> loc(#loc157)
+    %tmp12 = arith.constant 0.000000e+00 : f32 loc(#loc158)
+    %tmp12_63 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc158)
+    %tmp12_64 = arith.subf %tmp12_63, %tmp11 : tensor<1024xf32> loc(#loc158)
+    %tmp13 = arith.constant 0.000000e+00 : f32 loc(#loc159)
+    %tmp13_65 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc159)
+    %tmp14 = arith.select %tmp2_19, %tmp12_64, %tmp13_65 : tensor<1024xi1>, tensor<1024xf32> loc(#loc160)
+    %tmp15 = arith.constant 0.000000e+00 : f32 loc(#loc161)
+    %tmp16 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc162)
+    %tmp16_66 = arith.select %tmp2_19, %tmp14, %tmp16 : tensor<1024xi1>, tensor<1024xf32> loc(#loc162)
+    %tmp17 = tt.splat %tmp1_18 : i64 -> tensor<1024xi64> loc(#loc163)
+    %tmp17_67 = arith.cmpi slt, %x0_7, %tmp17 : tensor<1024xi64> loc(#loc163)
+    %tmp18 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc164)
+    %tmp18_68 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc164)
+    %tmp18_69 = arith.addi %tmp18_68, %tmp18 : tensor<1024xi64> loc(#loc164)
+    %tmp18_70 = arith.constant 2 : i32 loc(#loc165)
+    %tmp18_71 = arith.constant 2 : i64 loc(#loc165)
+    %tmp18_72 = arith.divsi %ks0, %tmp18_71 : i64 loc(#loc165)
+    %tmp18_73 = arith.constant -1 : i32 loc(#loc166)
+    %tmp18_74 = arith.constant -1 : i64 loc(#loc166)
+    %tmp18_75 = arith.muli %tmp18_74, %tmp18_72 : i64 loc(#loc166)
+    %tmp18_76 = tt.splat %tmp18_75 : i64 -> tensor<1024xi64> loc(#loc167)
+    %tmp18_77 = arith.addi %tmp18_69, %tmp18_76 : tensor<1024xi64> loc(#loc167)
+    %tmp18_78 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc168)
+    %tmp18_79 = tt.addptr %tmp18_78, %tmp18_77 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc168)
+    %tmp18_80 = arith.andi %tmp17_67, %xmask_5 : tensor<1024xi1> loc(#loc169)
+    %tmp18_81 = arith.constant 0.000000e+00 : f32 loc(#loc170)
+    %tmp18_82 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc170)
+    %tmp18_83 = arith.truncf %tmp18_82 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc170)
+    %tmp18_84 = tt.load %tmp18_79, %tmp18_80, %tmp18_83 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc170)
+    %tmp18_85 = arith.extf %tmp18_84 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc171)
+    %tmp19 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>> loc(#loc172)
+    %tmp19_86 = tt.addptr %tmp19, %x1_11 : tensor<1024x!tt.ptr<i64>>, tensor<1024xi64> loc(#loc172)
+    %tmp19_87 = arith.andi %tmp17_67, %xmask_5 : tensor<1024xi1> loc(#loc173)
+    %tmp19_88 = arith.constant 0.000000e+00 : f32 loc(#loc174)
+    %tmp19_89 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc174)
+    %tmp19_90 = arith.fptosi %tmp19_89 : tensor<1024xf32> to tensor<1024xi64> loc(#loc174)
+    %tmp19_91 = tt.load %tmp19_86, %tmp19_87, %tmp19_90 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc174)
+    %tmp20 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc175)
+    %tmp21 = arith.addi %tmp19_91, %tmp20 : tensor<1024xi64> loc(#loc176)
+    %tmp22 = arith.constant 0 : i32 loc(#loc177)
+    %tmp22_92 = arith.extsi %tmp22 : i32 to i64 loc(#loc177)
+    %tmp22_93 = tt.splat %tmp22_92 : i64 -> tensor<1024xi64> loc(#loc177)
+    %tmp22_94 = arith.cmpi slt, %tmp19_91, %tmp22_93 : tensor<1024xi64> loc(#loc177)
+    %tmp23 = arith.select %tmp22_94, %tmp21, %tmp19_91 : tensor<1024xi1>, tensor<1024xi64> loc(#loc178)
+    %c0_i32_95 = arith.constant 0 : i32 loc(#loc68)
+    %9 = arith.extsi %c0_i32_95 : i32 to i64 loc(#loc68)
+    %10 = tt.splat %9 : i64 -> tensor<1024xi64> loc(#loc68)
+    %11 = arith.cmpi sle, %10, %tmp23 : tensor<1024xi64> loc(#loc68)
+    %12 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc69)
+    %13 = arith.cmpi slt, %tmp23, %12 : tensor<1024xi64> loc(#loc69)
+    %14 = arith.andi %11, %13 : tensor<1024xi1> loc(#loc70)
+    %15 = arith.andi %tmp17_67, %xmask_5 : tensor<1024xi1> loc(#loc71)
+    %true_96 = arith.constant true loc(#loc72)
+    %cst_97 = arith.constant dense<true> : tensor<1024xi1> loc(#loc72)
+    %16 = arith.xori %15, %cst_97 : tensor<1024xi1> loc(#loc72)
+    %17 = arith.ori %14, %16 : tensor<1024xi1> loc(#loc73)
+    tt.assert %17, "index out of bounds: 0 <= tl.broadcast_to(tmp23, [XBLOCK]) < ks2" : tensor<1024xi1> loc(#loc74)
+    %tmp25 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc179)
+    %tmp25_98 = arith.addi %tmp25, %x0_7 : tensor<1024xi64> loc(#loc179)
+    %tmp25_99 = arith.constant 2 : i32 loc(#loc180)
+    %tmp25_100 = arith.constant 2 : i64 loc(#loc180)
+    %tmp25_101 = arith.divsi %ks0, %tmp25_100 : i64 loc(#loc180)
+    %tmp25_102 = arith.constant -1 : i32 loc(#loc181)
+    %tmp25_103 = arith.constant -1 : i64 loc(#loc181)
+    %tmp25_104 = arith.muli %tmp25_103, %tmp25_101 : i64 loc(#loc181)
+    %tmp25_105 = tt.splat %tmp25_104 : i64 -> tensor<1024xi64> loc(#loc182)
+    %tmp25_106 = arith.addi %tmp25_98, %tmp25_105 : tensor<1024xi64> loc(#loc182)
+    %tmp25_107 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc183)
+    %tmp25_108 = arith.muli %tmp25_107, %tmp23 : tensor<1024xi64> loc(#loc183)
+    %tmp25_109 = arith.addi %tmp25_106, %tmp25_108 : tensor<1024xi64> loc(#loc184)
+    %tmp25_110 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc185)
+    %tmp25_111 = tt.addptr %tmp25_110, %tmp25_109 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc185)
+    %tmp25_112 = arith.andi %tmp17_67, %xmask_5 : tensor<1024xi1> loc(#loc186)
+    %tmp25_113 = arith.constant 0.000000e+00 : f32 loc(#loc187)
+    %tmp25_114 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc187)
+    %tmp25_115 = arith.truncf %tmp25_114 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc187)
+    %tmp25_116 = tt.load %tmp25_111, %tmp25_112, %tmp25_115 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc187)
+    %tmp25_117 = arith.extf %tmp25_116 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc188)
+    %tmp26 = arith.mulf %tmp18_85, %tmp25_117 : tensor<1024xf32> loc(#loc189)
+    %tmp27 = arith.constant 0.000000e+00 : f32 loc(#loc190)
+    %tmp27_118 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc190)
+    %tmp28 = arith.select %tmp17_67, %tmp26, %tmp27_118 : tensor<1024xi1>, tensor<1024xf32> loc(#loc191)
+    %tmp29 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc192)
+    %tmp29_119 = arith.select %tmp17_67, %tmp28, %tmp29 : tensor<1024xi1>, tensor<1024xf32> loc(#loc192)
+    %tmp30 = arith.addf %tmp16_66, %tmp29_119 : tensor<1024xf32> loc(#loc193)
+    %tmp34 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc194)
+    %tmp34_120 = arith.addi %tmp32_16, %tmp34 : tensor<1024xi64> loc(#loc194)
+    %tmp35 = arith.constant 0 : i32 loc(#loc195)
+    %tmp35_121 = arith.extsi %tmp35 : i32 to i64 loc(#loc195)
+    %tmp35_122 = tt.splat %tmp35_121 : i64 -> tensor<1024xi64> loc(#loc195)
+    %tmp35_123 = arith.cmpi slt, %tmp32_16, %tmp35_122 : tensor<1024xi64> loc(#loc195)
+    %tmp36 = arith.select %tmp35_123, %tmp34_120, %tmp32_16 : tensor<1024xi1>, tensor<1024xi64> loc(#loc196)
+    %c0_i32_124 = arith.constant 0 : i32 loc(#loc93)
+    %18 = arith.extsi %c0_i32_124 : i32 to i64 loc(#loc93)
+    %19 = tt.splat %18 : i64 -> tensor<1024xi64> loc(#loc93)
+    %20 = arith.cmpi sle, %19, %tmp36 : tensor<1024xi64> loc(#loc93)
+    %21 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc94)
+    %22 = arith.cmpi slt, %tmp36, %21 : tensor<1024xi64> loc(#loc94)
+    %23 = arith.andi %20, %22 : tensor<1024xi1> loc(#loc95)
+    %true_125 = arith.constant true loc(#loc96)
+    %cst_126 = arith.constant dense<true> : tensor<1024xi1> loc(#loc96)
+    %24 = arith.xori %xmask_5, %cst_126 : tensor<1024xi1> loc(#loc96)
+    %25 = arith.ori %23, %24 : tensor<1024xi1> loc(#loc97)
+    tt.assert %25, "index out of bounds: 0 <= tmp36 < ks3" : tensor<1024xi1> loc(#loc98)
+    %tmp38 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc197)
+    %tmp38_127 = arith.muli %tmp38, %tmp36 : tensor<1024xi64> loc(#loc197)
+    %tmp38_128 = arith.addi %x0_7, %tmp38_127 : tensor<1024xi64> loc(#loc198)
+    %tmp38_129 = tt.splat %in_ptr3 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc199)
+    %tmp38_130 = tt.addptr %tmp38_129, %tmp38_128 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc199)
+    %tmp38_131 = tt.load %tmp38_130, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc200)
+    %tmp38_132 = arith.extf %tmp38_131 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc201)
+    %tmp39 = arith.mulf %tmp31_14, %tmp38_132 : tensor<1024xf32> loc(#loc202)
+    %tmp40 = arith.addf %tmp30, %tmp39 : tensor<1024xf32> loc(#loc203)
+    %26 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc106)
+    %27 = tt.addptr %26, %xindex_4 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc106)
+    %28 = arith.truncf %tmp40 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc107)
+    tt.store %27, %28, %xmask_5 : tensor<1024x!tt.ptr<bf16>> loc(#loc107)
+    tt.return loc(#loc108)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":19:28)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":19:33)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":20:36)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":20:23)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":21:21)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":22:19)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":24:21)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":24:28)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:31)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:36)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:76)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":26:31)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":26:36)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":28:18)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":29:19)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:48)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:41)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:35)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:30)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:60)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:53)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:111)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":31:30)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":31:42)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":31:35)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":32:32)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":33:18)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":34:18)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":35:32)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:28)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:98)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:64)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:115)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:108)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:106)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:123)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:49)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:42)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:36)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:58)
+#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:54)
+#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:31)
+#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:72)
+#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:65)
+#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:123)
+#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":38:19)
+#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":39:13)
+#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":40:38)
+#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":41:34)
+#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":42:12)
+#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":43:34)
+#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":44:19)
+#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:37)
+#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:55)
+#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:48)
+#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:42)
+#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:31)
+#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:68)
+#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:60)
+#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:119)
+#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":46:31)
+#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":46:44)
+#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":46:36)
+#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":47:33)
+#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":48:20)
+#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":49:20)
+#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":50:35)
+#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:28)
+#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:100)
+#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:65)
+#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:118)
+#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:110)
+#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:108)
+#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:126)
+#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:37)
+#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:55)
+#loc77 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:48)
+#loc78 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:42)
+#loc79 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:64)
+#loc80 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:60)
+#loc81 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:31)
+#loc82 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:80)
+#loc83 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:72)
+#loc84 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:131)
+#loc85 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":53:20)
+#loc86 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":54:38)
+#loc87 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":55:35)
+#loc88 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":56:35)
+#loc89 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":57:20)
+#loc90 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":59:20)
+#loc91 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":60:20)
+#loc92 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":61:35)
+#loc93 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:28)
+#loc94 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:46)
+#loc95 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:38)
+#loc96 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:56)
+#loc97 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:54)
+#loc98 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:64)
+#loc99 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:40)
+#loc100 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:36)
+#loc101 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:31)
+#loc102 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:48)
+#loc103 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:88)
+#loc104 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":64:20)
+#loc105 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":65:20)
+#loc106 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:25)
+#loc107 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:37)
+#loc108 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:4)
+#loc119 = loc("xoffset"(#loc1))
+#loc120 = loc("xoffset"(#loc2))
+#loc121 = loc("xindex"(#loc3))
+#loc122 = loc("xindex"(#loc4))
+#loc123 = loc("xmask"(#loc5))
+#loc124 = loc("x0"(#loc6))
+#loc125 = loc("x1"(#loc7))
+#loc126 = loc("x1"(#loc8))
+#loc127 = loc("tmp31"(#loc9))
+#loc128 = loc("tmp31"(#loc10))
+#loc129 = loc("tmp31"(#loc11))
+#loc130 = loc("tmp32"(#loc12))
+#loc131 = loc("tmp32"(#loc13))
+#loc132 = loc("tmp1"(#loc14))
+#loc133 = loc("tmp2"(#loc15))
+#loc134 = loc("tmp3"(#loc16))
+#loc135 = loc("tmp3"(#loc17))
+#loc136 = loc("tmp3"(#loc18))
+#loc137 = loc("tmp3"(#loc19))
+#loc138 = loc("tmp3"(#loc20))
+#loc139 = loc("tmp3"(#loc21))
+#loc140 = loc("tmp3"(#loc22))
+#loc141 = loc("tmp4"(#loc23))
+#loc142 = loc("tmp4"(#loc24))
+#loc143 = loc("tmp4"(#loc25))
+#loc144 = loc("tmp5"(#loc26))
+#loc145 = loc("tmp6"(#loc27))
+#loc146 = loc("tmp7"(#loc28))
+#loc147 = loc("tmp8"(#loc29))
+#loc148 = loc("tmp10"(#loc37))
+#loc149 = loc("tmp10"(#loc38))
+#loc150 = loc("tmp10"(#loc39))
+#loc151 = loc("tmp10"(#loc40))
+#loc152 = loc("tmp10"(#loc41))
+#loc153 = loc("tmp10"(#loc42))
+#loc154 = loc("tmp10"(#loc43))
+#loc155 = loc("tmp10"(#loc44))
+#loc156 = loc("tmp10"(#loc45))
+#loc157 = loc("tmp11"(#loc46))
+#loc158 = loc("tmp12"(#loc47))
+#loc159 = loc("tmp13"(#loc48))
+#loc160 = loc("tmp14"(#loc49))
+#loc161 = loc("tmp15"(#loc50))
+#loc162 = loc("tmp16"(#loc51))
+#loc163 = loc("tmp17"(#loc52))
+#loc164 = loc("tmp18"(#loc53))
+#loc165 = loc("tmp18"(#loc54))
+#loc166 = loc("tmp18"(#loc55))
+#loc167 = loc("tmp18"(#loc56))
+#loc168 = loc("tmp18"(#loc57))
+#loc169 = loc("tmp18"(#loc58))
+#loc170 = loc("tmp18"(#loc59))
+#loc171 = loc("tmp18"(#loc60))
+#loc172 = loc("tmp19"(#loc61))
+#loc173 = loc("tmp19"(#loc62))
+#loc174 = loc("tmp19"(#loc63))
+#loc175 = loc("tmp20"(#loc64))
+#loc176 = loc("tmp21"(#loc65))
+#loc177 = loc("tmp22"(#loc66))
+#loc178 = loc("tmp23"(#loc67))
+#loc179 = loc("tmp25"(#loc75))
+#loc180 = loc("tmp25"(#loc76))
+#loc181 = loc("tmp25"(#loc77))
+#loc182 = loc("tmp25"(#loc78))
+#loc183 = loc("tmp25"(#loc79))
+#loc184 = loc("tmp25"(#loc80))
+#loc185 = loc("tmp25"(#loc81))
+#loc186 = loc("tmp25"(#loc82))
+#loc187 = loc("tmp25"(#loc83))
+#loc188 = loc("tmp25"(#loc84))
+#loc189 = loc("tmp26"(#loc85))
+#loc190 = loc("tmp27"(#loc86))
+#loc191 = loc("tmp28"(#loc87))
+#loc192 = loc("tmp29"(#loc88))
+#loc193 = loc("tmp30"(#loc89))
+#loc194 = loc("tmp34"(#loc90))
+#loc195 = loc("tmp35"(#loc91))
+#loc196 = loc("tmp36"(#loc92))
+#loc197 = loc("tmp38"(#loc99))
+#loc198 = loc("tmp38"(#loc100))
+#loc199 = loc("tmp38"(#loc101))
+#loc200 = loc("tmp38"(#loc102))
+#loc201 = loc("tmp38"(#loc103))
+#loc202 = loc("tmp39"(#loc104))
+#loc203 = loc("tmp40"(#loc105))

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ttgir

@@ -0,0 +1,284 @@
+#blocked = #ttg.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":18:0)
+#loc91 = loc("in_ptr0"(#loc))
+#loc92 = loc("in_ptr1"(#loc))
+#loc93 = loc("in_ptr2"(#loc))
+#loc94 = loc("in_ptr3"(#loc))
+#loc95 = loc("out_ptr0"(#loc))
+#loc96 = loc("ks0"(#loc))
+#loc97 = loc("ks1"(#loc))
+#loc98 = loc("ks2"(#loc))
+#loc99 = loc("ks3"(#loc))
+#loc100 = loc("xnumel"(#loc))
+module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
+  tt.func public @triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %in_ptr2: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr2"(#loc)), %in_ptr3: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr3"(#loc)), %out_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %ks2: i64 loc("ks2"(#loc)), %ks3: i64 loc("ks3"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<true> : tensor<1024xi1, #blocked> loc(#loc1)
+    %c-1_i64 = arith.constant -1 : i64 loc(#loc1)
+    %c2_i64 = arith.constant 2 : i64 loc(#loc1)
+    %c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
+    %cst_0 = arith.constant dense<0.000000e+00> : tensor<1024xbf16, #blocked> loc(#loc1)
+    %cst_1 = arith.constant dense<0> : tensor<1024xi64, #blocked> loc(#loc1)
+    %cst_2 = arith.constant dense<0.000000e+00> : tensor<1024xf32, #blocked> loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc101)
+    %xoffset_3 = arith.muli %xoffset, %c1024_i32 : i32 loc(#loc102)
+    %xindex = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #blocked> loc(#loc103)
+    %xindex_4 = tt.splat %xoffset_3 : i32 -> tensor<1024xi32, #blocked> loc(#loc104)
+    %xindex_5 = arith.addi %xindex_4, %xindex : tensor<1024xi32, #blocked> loc(#loc104)
+    %xmask = tt.splat %xnumel : i32 -> tensor<1024xi32, #blocked> loc(#loc105)
+    %xmask_6 = arith.cmpi slt, %xindex_5, %xmask : tensor<1024xi32, #blocked> loc(#loc105)
+    %x0 = arith.extsi %xindex_5 : tensor<1024xi32, #blocked> to tensor<1024xi64, #blocked> loc(#loc106)
+    %x0_7 = tt.splat %ks0 : i64 -> tensor<1024xi64, #blocked> loc(#loc106)
+    %x0_8 = arith.remsi %x0, %x0_7 : tensor<1024xi64, #blocked> loc(#loc106)
+    %x1 = arith.divsi %x0, %x0_7 : tensor<1024xi64, #blocked> loc(#loc107)
+    %x1_9 = tt.splat %ks1 : i64 -> tensor<1024xi64, #blocked> loc(#loc108)
+    %x1_10 = arith.remsi %x1, %x1_9 : tensor<1024xi64, #blocked> loc(#loc108)
+    %tmp31 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc109)
+    %tmp31_11 = tt.addptr %tmp31, %xindex_5 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi32, #blocked> loc(#loc109)
+    %tmp31_12 = tt.load %tmp31_11, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc110)
+    %tmp31_13 = arith.extf %tmp31_12 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc111)
+    %tmp32 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>, #blocked> loc(#loc112)
+    %tmp32_14 = tt.addptr %tmp32, %x1_10 : tensor<1024x!tt.ptr<i64>, #blocked>, tensor<1024xi64, #blocked> loc(#loc112)
+    %tmp32_15 = tt.load %tmp32_14, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>, #blocked> loc(#loc113)
+    %tmp1 = arith.divsi %ks0, %c2_i64 : i64 loc(#loc114)
+    %tmp2 = tt.splat %tmp1 : i64 -> tensor<1024xi64, #blocked> loc(#loc115)
+    %tmp2_16 = arith.cmpi sge, %x0_8, %tmp2 : tensor<1024xi64, #blocked> loc(#loc115)
+    %tmp3 = arith.muli %tmp1, %c-1_i64 : i64 loc(#loc116)
+    %tmp3_17 = tt.splat %tmp3 : i64 -> tensor<1024xi64, #blocked> loc(#loc117)
+    %tmp3_18 = arith.addi %x0, %tmp3_17 : tensor<1024xi64, #blocked> loc(#loc117)
+    %tmp3_19 = tt.addptr %tmp31, %tmp3_18 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc118)
+    %tmp3_20 = arith.andi %tmp2_16, %xmask_6 : tensor<1024xi1, #blocked> loc(#loc119)
+    %tmp3_21 = tt.load %tmp3_19, %tmp3_20, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc120)
+    %tmp3_22 = arith.extf %tmp3_21 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc121)
+    %tmp4 = tt.load %tmp32_14, %tmp3_20, %cst_1 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>, #blocked> loc(#loc122)
+    %tmp5 = tt.splat %ks2 : i64 -> tensor<1024xi64, #blocked> loc(#loc123)
+    %tmp6 = arith.addi %tmp4, %tmp5 : tensor<1024xi64, #blocked> loc(#loc124)
+    %tmp7 = arith.cmpi slt, %tmp4, %cst_1 : tensor<1024xi64, #blocked> loc(#loc125)
+    %tmp8 = arith.select %tmp7, %tmp6, %tmp4 : tensor<1024xi1, #blocked>, tensor<1024xi64, #blocked> loc(#loc126)
+    %0 = arith.cmpi sge, %tmp8, %cst_1 : tensor<1024xi64, #blocked> loc(#loc28)
+    %1 = arith.cmpi slt, %tmp8, %tmp5 : tensor<1024xi64, #blocked> loc(#loc29)
+    %2 = arith.andi %0, %1 : tensor<1024xi1, #blocked> loc(#loc30)
+    %3 = arith.xori %tmp3_20, %cst : tensor<1024xi1, #blocked> loc(#loc31)
+    %4 = arith.ori %2, %3 : tensor<1024xi1, #blocked> loc(#loc32)
+    tt.assert %4, "index out of bounds: 0 <= tl.broadcast_to(tmp8, [XBLOCK]) < ks2" : tensor<1024xi1, #blocked> loc(#loc33)
+    %tmp10 = arith.addi %x0_8, %tmp3_17 : tensor<1024xi64, #blocked> loc(#loc127)
+    %tmp10_23 = arith.muli %x0_7, %tmp8 : tensor<1024xi64, #blocked> loc(#loc128)
+    %tmp10_24 = arith.addi %tmp10, %tmp10_23 : tensor<1024xi64, #blocked> loc(#loc129)
+    %tmp10_25 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc130)
+    %tmp10_26 = tt.addptr %tmp10_25, %tmp10_24 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc130)
+    %tmp10_27 = tt.load %tmp10_26, %tmp3_20, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc131)
+    %tmp10_28 = arith.extf %tmp10_27 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc132)
+    %tmp11 = arith.mulf %tmp3_22, %tmp10_28 : tensor<1024xf32, #blocked> loc(#loc133)
+    %tmp12 = arith.subf %cst_2, %tmp11 : tensor<1024xf32, #blocked> loc(#loc134)
+    %tmp16 = arith.select %tmp2_16, %tmp12, %cst_2 : tensor<1024xi1, #blocked>, tensor<1024xf32, #blocked> loc(#loc169)
+    %tmp17 = arith.cmpi slt, %x0_8, %tmp2 : tensor<1024xi64, #blocked> loc(#loc137)
+    %tmp18 = arith.addi %x0_7, %x0 : tensor<1024xi64, #blocked> loc(#loc138)
+    %tmp18_29 = arith.addi %tmp18, %tmp3_17 : tensor<1024xi64, #blocked> loc(#loc139)
+    %tmp18_30 = tt.addptr %tmp31, %tmp18_29 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc140)
+    %tmp18_31 = arith.andi %tmp17, %xmask_6 : tensor<1024xi1, #blocked> loc(#loc141)
+    %tmp18_32 = tt.load %tmp18_30, %tmp18_31, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc142)
+    %tmp18_33 = arith.extf %tmp18_32 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc143)
+    %tmp19 = tt.load %tmp32_14, %tmp18_31, %cst_1 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>, #blocked> loc(#loc144)
+    %tmp21 = arith.addi %tmp19, %tmp5 : tensor<1024xi64, #blocked> loc(#loc145)
+    %tmp22 = arith.cmpi slt, %tmp19, %cst_1 : tensor<1024xi64, #blocked> loc(#loc146)
+    %tmp23 = arith.select %tmp22, %tmp21, %tmp19 : tensor<1024xi1, #blocked>, tensor<1024xi64, #blocked> loc(#loc147)
+    %5 = arith.cmpi sge, %tmp23, %cst_1 : tensor<1024xi64, #blocked> loc(#loc55)
+    %6 = arith.cmpi slt, %tmp23, %tmp5 : tensor<1024xi64, #blocked> loc(#loc56)
+    %7 = arith.andi %5, %6 : tensor<1024xi1, #blocked> loc(#loc57)
+    %8 = arith.xori %tmp18_31, %cst : tensor<1024xi1, #blocked> loc(#loc58)
+    %9 = arith.ori %7, %8 : tensor<1024xi1, #blocked> loc(#loc59)
+    tt.assert %9, "index out of bounds: 0 <= tl.broadcast_to(tmp23, [XBLOCK]) < ks2" : tensor<1024xi1, #blocked> loc(#loc60)
+    %tmp25 = arith.addi %x0_7, %x0_8 : tensor<1024xi64, #blocked> loc(#loc148)
+    %tmp25_34 = arith.addi %tmp25, %tmp3_17 : tensor<1024xi64, #blocked> loc(#loc149)
+    %tmp25_35 = arith.muli %x0_7, %tmp23 : tensor<1024xi64, #blocked> loc(#loc150)
+    %tmp25_36 = arith.addi %tmp25_34, %tmp25_35 : tensor<1024xi64, #blocked> loc(#loc151)
+    %tmp25_37 = tt.addptr %tmp10_25, %tmp25_36 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc152)
+    %tmp25_38 = tt.load %tmp25_37, %tmp18_31, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc153)
+    %tmp25_39 = arith.extf %tmp25_38 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc154)
+    %tmp26 = arith.mulf %tmp18_33, %tmp25_39 : tensor<1024xf32, #blocked> loc(#loc155)
+    %tmp29 = arith.select %tmp17, %tmp26, %cst_2 : tensor<1024xi1, #blocked>, tensor<1024xf32, #blocked> loc(#loc170)
+    %tmp30 = arith.addf %tmp16, %tmp29 : tensor<1024xf32, #blocked> loc(#loc158)
+    %tmp34 = tt.splat %ks3 : i64 -> tensor<1024xi64, #blocked> loc(#loc159)
+    %tmp34_40 = arith.addi %tmp32_15, %tmp34 : tensor<1024xi64, #blocked> loc(#loc159)
+    %tmp35 = arith.cmpi slt, %tmp32_15, %cst_1 : tensor<1024xi64, #blocked> loc(#loc160)
+    %tmp36 = arith.select %tmp35, %tmp34_40, %tmp32_15 : tensor<1024xi1, #blocked>, tensor<1024xi64, #blocked> loc(#loc161)
+    %10 = arith.cmpi sge, %tmp36, %cst_1 : tensor<1024xi64, #blocked> loc(#loc75)
+    %11 = arith.cmpi slt, %tmp36, %tmp34 : tensor<1024xi64, #blocked> loc(#loc76)
+    %12 = arith.andi %10, %11 : tensor<1024xi1, #blocked> loc(#loc77)
+    %13 = arith.xori %xmask_6, %cst : tensor<1024xi1, #blocked> loc(#loc78)
+    %14 = arith.ori %12, %13 : tensor<1024xi1, #blocked> loc(#loc79)
+    tt.assert %14, "index out of bounds: 0 <= tmp36 < ks3" : tensor<1024xi1, #blocked> loc(#loc80)
+    %tmp38 = arith.muli %x0_7, %tmp36 : tensor<1024xi64, #blocked> loc(#loc162)
+    %tmp38_41 = arith.addi %x0_8, %tmp38 : tensor<1024xi64, #blocked> loc(#loc163)
+    %tmp38_42 = tt.splat %in_ptr3 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc164)
+    %tmp38_43 = tt.addptr %tmp38_42, %tmp38_41 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc164)
+    %tmp38_44 = tt.load %tmp38_43, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc165)
+    %tmp38_45 = arith.extf %tmp38_44 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc166)
+    %tmp39 = arith.mulf %tmp31_13, %tmp38_45 : tensor<1024xf32, #blocked> loc(#loc167)
+    %tmp40 = arith.addf %tmp30, %tmp39 : tensor<1024xf32, #blocked> loc(#loc168)
+    %15 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc88)
+    %16 = tt.addptr %15, %xindex_5 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi32, #blocked> loc(#loc88)
+    %17 = arith.truncf %tmp40 : tensor<1024xf32, #blocked> to tensor<1024xbf16, #blocked> loc(#loc89)
+    tt.store %16, %17, %xmask_6 : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc89)
+    tt.return loc(#loc90)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc(unknown)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":19:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":19:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":20:36)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":20:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":21:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":22:19)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":24:21)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":24:28)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:31)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:36)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:76)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":26:31)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":26:36)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":28:18)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":29:19)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:41)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:35)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:30)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:60)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:53)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:111)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":31:35)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":32:32)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":33:18)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":34:18)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":35:32)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:28)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:98)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:64)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:108)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:106)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:123)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:36)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:58)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:54)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:31)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:65)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:123)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":38:19)
+#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":39:13)
+#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":43:34)
+#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":41:34)
+#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":44:19)
+#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:37)
+#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:42)
+#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:31)
+#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:68)
+#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:60)
+#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:119)
+#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":46:36)
+#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":48:20)
+#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":49:20)
+#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":50:35)
+#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:28)
+#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:100)
+#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:65)
+#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:110)
+#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:108)
+#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:126)
+#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:37)
+#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:42)
+#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:64)
+#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:60)
+#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:31)
+#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:72)
+#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:131)
+#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":53:20)
+#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":56:35)
+#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":55:35)
+#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":57:20)
+#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":59:20)
+#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":60:20)
+#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":61:35)
+#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:28)
+#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:46)
+#loc77 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:38)
+#loc78 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:56)
+#loc79 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:54)
+#loc80 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:64)
+#loc81 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:40)
+#loc82 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:36)
+#loc83 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:31)
+#loc84 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:48)
+#loc85 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:88)
+#loc86 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":64:20)
+#loc87 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":65:20)
+#loc88 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:25)
+#loc89 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:37)
+#loc90 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:4)
+#loc101 = loc("xoffset"(#loc2))
+#loc102 = loc("xoffset"(#loc3))
+#loc103 = loc("xindex"(#loc4))
+#loc104 = loc("xindex"(#loc5))
+#loc105 = loc("xmask"(#loc6))
+#loc106 = loc("x0"(#loc7))
+#loc107 = loc("x1"(#loc8))
+#loc108 = loc("x1"(#loc9))
+#loc109 = loc("tmp31"(#loc10))
+#loc110 = loc("tmp31"(#loc11))
+#loc111 = loc("tmp31"(#loc12))
+#loc112 = loc("tmp32"(#loc13))
+#loc113 = loc("tmp32"(#loc14))
+#loc114 = loc("tmp1"(#loc15))
+#loc115 = loc("tmp2"(#loc16))
+#loc116 = loc("tmp3"(#loc17))
+#loc117 = loc("tmp3"(#loc18))
+#loc118 = loc("tmp3"(#loc19))
+#loc119 = loc("tmp3"(#loc20))
+#loc120 = loc("tmp3"(#loc21))
+#loc121 = loc("tmp3"(#loc22))
+#loc122 = loc("tmp4"(#loc23))
+#loc123 = loc("tmp5"(#loc24))
+#loc124 = loc("tmp6"(#loc25))
+#loc125 = loc("tmp7"(#loc26))
+#loc126 = loc("tmp8"(#loc27))
+#loc127 = loc("tmp10"(#loc34))
+#loc128 = loc("tmp10"(#loc35))
+#loc129 = loc("tmp10"(#loc36))
+#loc130 = loc("tmp10"(#loc37))
+#loc131 = loc("tmp10"(#loc38))
+#loc132 = loc("tmp10"(#loc39))
+#loc133 = loc("tmp11"(#loc40))
+#loc134 = loc("tmp12"(#loc41))
+#loc135 = loc("tmp16"(#loc42))
+#loc136 = loc("tmp14"(#loc43))
+#loc137 = loc("tmp17"(#loc44))
+#loc138 = loc("tmp18"(#loc45))
+#loc139 = loc("tmp18"(#loc46))
+#loc140 = loc("tmp18"(#loc47))
+#loc141 = loc("tmp18"(#loc48))
+#loc142 = loc("tmp18"(#loc49))
+#loc143 = loc("tmp18"(#loc50))
+#loc144 = loc("tmp19"(#loc51))
+#loc145 = loc("tmp21"(#loc52))
+#loc146 = loc("tmp22"(#loc53))
+#loc147 = loc("tmp23"(#loc54))
+#loc148 = loc("tmp25"(#loc61))
+#loc149 = loc("tmp25"(#loc62))
+#loc150 = loc("tmp25"(#loc63))
+#loc151 = loc("tmp25"(#loc64))
+#loc152 = loc("tmp25"(#loc65))
+#loc153 = loc("tmp25"(#loc66))
+#loc154 = loc("tmp25"(#loc67))
+#loc155 = loc("tmp26"(#loc68))
+#loc156 = loc("tmp29"(#loc69))
+#loc157 = loc("tmp28"(#loc70))
+#loc158 = loc("tmp30"(#loc71))
+#loc159 = loc("tmp34"(#loc72))
+#loc160 = loc("tmp35"(#loc73))
+#loc161 = loc("tmp36"(#loc74))
+#loc162 = loc("tmp38"(#loc81))
+#loc163 = loc("tmp38"(#loc82))
+#loc164 = loc("tmp38"(#loc83))
+#loc165 = loc("tmp38"(#loc84))
+#loc166 = loc("tmp38"(#loc85))
+#loc167 = loc("tmp39"(#loc86))
+#loc168 = loc("tmp40"(#loc87))
+#loc169 = loc(fused[#loc135, #loc136])
+#loc170 = loc(fused[#loc156, #loc157])

SpecForge-ext/cache/compiled_kernels/triton/6/FWEPGR3FZNND66AYQXX4GQT7TCIT4B6L5XAV2CXTIKJFUO4ZTCVQ/triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0.ttir

@@ -0,0 +1,283 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":18:0)
+#loc91 = loc("in_ptr0"(#loc))
+#loc92 = loc("in_ptr1"(#loc))
+#loc93 = loc("in_ptr2"(#loc))
+#loc94 = loc("in_ptr3"(#loc))
+#loc95 = loc("out_ptr0"(#loc))
+#loc96 = loc("ks0"(#loc))
+#loc97 = loc("ks1"(#loc))
+#loc98 = loc("ks2"(#loc))
+#loc99 = loc("ks3"(#loc))
+#loc100 = loc("xnumel"(#loc))
+module {
+  tt.func public @triton_poi_fused_add_index_mul_neg_slice_slice_backward_squeeze_unsqueeze_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %in_ptr2: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr2"(#loc)), %in_ptr3: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr3"(#loc)), %out_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %ks2: i64 loc("ks2"(#loc)), %ks3: i64 loc("ks3"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<0> : tensor<1024xi64> loc(#loc1)
+    %cst_0 = arith.constant dense<0.000000e+00> : tensor<1024xbf16> loc(#loc1)
+    %cst_1 = arith.constant dense<true> : tensor<1024xi1> loc(#loc1)
+    %cst_2 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc1)
+    %c-1_i64 = arith.constant -1 : i64 loc(#loc1)
+    %c2_i64 = arith.constant 2 : i64 loc(#loc1)
+    %c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc101)
+    %xoffset_3 = arith.muli %xoffset, %c1024_i32 : i32 loc(#loc102)
+    %xindex = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc103)
+    %xindex_4 = tt.splat %xoffset_3 : i32 -> tensor<1024xi32> loc(#loc104)
+    %xindex_5 = arith.addi %xindex_4, %xindex : tensor<1024xi32> loc(#loc104)
+    %xmask = tt.splat %xnumel : i32 -> tensor<1024xi32> loc(#loc105)
+    %xmask_6 = arith.cmpi slt, %xindex_5, %xmask : tensor<1024xi32> loc(#loc105)
+    %x0 = arith.extsi %xindex_5 : tensor<1024xi32> to tensor<1024xi64> loc(#loc106)
+    %x0_7 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc106)
+    %x0_8 = arith.remsi %x0, %x0_7 : tensor<1024xi64> loc(#loc106)
+    %x1 = arith.divsi %x0, %x0_7 : tensor<1024xi64> loc(#loc107)
+    %x1_9 = tt.splat %ks1 : i64 -> tensor<1024xi64> loc(#loc108)
+    %x1_10 = arith.remsi %x1, %x1_9 : tensor<1024xi64> loc(#loc108)
+    %tmp31 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc109)
+    %tmp31_11 = tt.addptr %tmp31, %xindex_5 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc109)
+    %tmp31_12 = tt.load %tmp31_11, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc110)
+    %tmp31_13 = arith.extf %tmp31_12 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc111)
+    %tmp32 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>> loc(#loc112)
+    %tmp32_14 = tt.addptr %tmp32, %x1_10 : tensor<1024x!tt.ptr<i64>>, tensor<1024xi64> loc(#loc112)
+    %tmp32_15 = tt.load %tmp32_14, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc113)
+    %tmp1 = arith.divsi %ks0, %c2_i64 : i64 loc(#loc114)
+    %tmp2 = tt.splat %tmp1 : i64 -> tensor<1024xi64> loc(#loc115)
+    %tmp2_16 = arith.cmpi sge, %x0_8, %tmp2 : tensor<1024xi64> loc(#loc115)
+    %tmp3 = arith.muli %tmp1, %c-1_i64 : i64 loc(#loc116)
+    %tmp3_17 = tt.splat %tmp3 : i64 -> tensor<1024xi64> loc(#loc117)
+    %tmp3_18 = arith.addi %x0, %tmp3_17 : tensor<1024xi64> loc(#loc117)
+    %tmp3_19 = tt.addptr %tmp31, %tmp3_18 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc118)
+    %tmp3_20 = arith.andi %tmp2_16, %xmask_6 : tensor<1024xi1> loc(#loc119)
+    %tmp3_21 = tt.load %tmp3_19, %tmp3_20, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc120)
+    %tmp3_22 = arith.extf %tmp3_21 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc121)
+    %tmp4 = tt.load %tmp32_14, %tmp3_20, %cst evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc122)
+    %tmp5 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc123)
+    %tmp6 = arith.addi %tmp4, %tmp5 : tensor<1024xi64> loc(#loc124)
+    %tmp7 = arith.cmpi slt, %tmp4, %cst : tensor<1024xi64> loc(#loc125)
+    %tmp8 = arith.select %tmp7, %tmp6, %tmp4 : tensor<1024xi1>, tensor<1024xi64> loc(#loc126)
+    %0 = arith.cmpi sge, %tmp8, %cst : tensor<1024xi64> loc(#loc28)
+    %1 = arith.cmpi slt, %tmp8, %tmp5 : tensor<1024xi64> loc(#loc29)
+    %2 = arith.andi %0, %1 : tensor<1024xi1> loc(#loc30)
+    %3 = arith.xori %tmp3_20, %cst_1 : tensor<1024xi1> loc(#loc31)
+    %4 = arith.ori %2, %3 : tensor<1024xi1> loc(#loc32)
+    tt.assert %4, "index out of bounds: 0 <= tl.broadcast_to(tmp8, [XBLOCK]) < ks2" : tensor<1024xi1> loc(#loc33)
+    %tmp10 = arith.addi %x0_8, %tmp3_17 : tensor<1024xi64> loc(#loc127)
+    %tmp10_23 = arith.muli %x0_7, %tmp8 : tensor<1024xi64> loc(#loc128)
+    %tmp10_24 = arith.addi %tmp10, %tmp10_23 : tensor<1024xi64> loc(#loc129)
+    %tmp10_25 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc130)
+    %tmp10_26 = tt.addptr %tmp10_25, %tmp10_24 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc130)
+    %tmp10_27 = tt.load %tmp10_26, %tmp3_20, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc131)
+    %tmp10_28 = arith.extf %tmp10_27 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc132)
+    %tmp11 = arith.mulf %tmp3_22, %tmp10_28 : tensor<1024xf32> loc(#loc133)
+    %tmp12 = arith.subf %cst_2, %tmp11 : tensor<1024xf32> loc(#loc134)
+    %tmp16 = arith.select %tmp2_16, %tmp12, %cst_2 : tensor<1024xi1>, tensor<1024xf32> loc(#loc169)
+    %tmp17 = arith.cmpi slt, %x0_8, %tmp2 : tensor<1024xi64> loc(#loc137)
+    %tmp18 = arith.addi %x0_7, %x0 : tensor<1024xi64> loc(#loc138)
+    %tmp18_29 = arith.addi %tmp18, %tmp3_17 : tensor<1024xi64> loc(#loc139)
+    %tmp18_30 = tt.addptr %tmp31, %tmp18_29 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc140)
+    %tmp18_31 = arith.andi %tmp17, %xmask_6 : tensor<1024xi1> loc(#loc141)
+    %tmp18_32 = tt.load %tmp18_30, %tmp18_31, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc142)
+    %tmp18_33 = arith.extf %tmp18_32 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc143)
+    %tmp19 = tt.load %tmp32_14, %tmp18_31, %cst evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc144)
+    %tmp21 = arith.addi %tmp19, %tmp5 : tensor<1024xi64> loc(#loc145)
+    %tmp22 = arith.cmpi slt, %tmp19, %cst : tensor<1024xi64> loc(#loc146)
+    %tmp23 = arith.select %tmp22, %tmp21, %tmp19 : tensor<1024xi1>, tensor<1024xi64> loc(#loc147)
+    %5 = arith.cmpi sge, %tmp23, %cst : tensor<1024xi64> loc(#loc55)
+    %6 = arith.cmpi slt, %tmp23, %tmp5 : tensor<1024xi64> loc(#loc56)
+    %7 = arith.andi %5, %6 : tensor<1024xi1> loc(#loc57)
+    %8 = arith.xori %tmp18_31, %cst_1 : tensor<1024xi1> loc(#loc58)
+    %9 = arith.ori %7, %8 : tensor<1024xi1> loc(#loc59)
+    tt.assert %9, "index out of bounds: 0 <= tl.broadcast_to(tmp23, [XBLOCK]) < ks2" : tensor<1024xi1> loc(#loc60)
+    %tmp25 = arith.addi %x0_7, %x0_8 : tensor<1024xi64> loc(#loc148)
+    %tmp25_34 = arith.addi %tmp25, %tmp3_17 : tensor<1024xi64> loc(#loc149)
+    %tmp25_35 = arith.muli %x0_7, %tmp23 : tensor<1024xi64> loc(#loc150)
+    %tmp25_36 = arith.addi %tmp25_34, %tmp25_35 : tensor<1024xi64> loc(#loc151)
+    %tmp25_37 = tt.addptr %tmp10_25, %tmp25_36 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc152)
+    %tmp25_38 = tt.load %tmp25_37, %tmp18_31, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc153)
+    %tmp25_39 = arith.extf %tmp25_38 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc154)
+    %tmp26 = arith.mulf %tmp18_33, %tmp25_39 : tensor<1024xf32> loc(#loc155)
+    %tmp29 = arith.select %tmp17, %tmp26, %cst_2 : tensor<1024xi1>, tensor<1024xf32> loc(#loc170)
+    %tmp30 = arith.addf %tmp16, %tmp29 : tensor<1024xf32> loc(#loc158)
+    %tmp34 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc159)
+    %tmp34_40 = arith.addi %tmp32_15, %tmp34 : tensor<1024xi64> loc(#loc159)
+    %tmp35 = arith.cmpi slt, %tmp32_15, %cst : tensor<1024xi64> loc(#loc160)
+    %tmp36 = arith.select %tmp35, %tmp34_40, %tmp32_15 : tensor<1024xi1>, tensor<1024xi64> loc(#loc161)
+    %10 = arith.cmpi sge, %tmp36, %cst : tensor<1024xi64> loc(#loc75)
+    %11 = arith.cmpi slt, %tmp36, %tmp34 : tensor<1024xi64> loc(#loc76)
+    %12 = arith.andi %10, %11 : tensor<1024xi1> loc(#loc77)
+    %13 = arith.xori %xmask_6, %cst_1 : tensor<1024xi1> loc(#loc78)
+    %14 = arith.ori %12, %13 : tensor<1024xi1> loc(#loc79)
+    tt.assert %14, "index out of bounds: 0 <= tmp36 < ks3" : tensor<1024xi1> loc(#loc80)
+    %tmp38 = arith.muli %x0_7, %tmp36 : tensor<1024xi64> loc(#loc162)
+    %tmp38_41 = arith.addi %x0_8, %tmp38 : tensor<1024xi64> loc(#loc163)
+    %tmp38_42 = tt.splat %in_ptr3 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc164)
+    %tmp38_43 = tt.addptr %tmp38_42, %tmp38_41 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc164)
+    %tmp38_44 = tt.load %tmp38_43, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc165)
+    %tmp38_45 = arith.extf %tmp38_44 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc166)
+    %tmp39 = arith.mulf %tmp31_13, %tmp38_45 : tensor<1024xf32> loc(#loc167)
+    %tmp40 = arith.addf %tmp30, %tmp39 : tensor<1024xf32> loc(#loc168)
+    %15 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc88)
+    %16 = tt.addptr %15, %xindex_5 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc88)
+    %17 = arith.truncf %tmp40 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc89)
+    tt.store %16, %17, %xmask_6 : tensor<1024x!tt.ptr<bf16>> loc(#loc89)
+    tt.return loc(#loc90)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc(unknown)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":19:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":19:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":20:36)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":20:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":21:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":22:19)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":24:21)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":24:28)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:31)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:36)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":25:76)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":26:31)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":26:36)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":28:18)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":29:19)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:41)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:35)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:30)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:60)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:53)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":30:111)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":31:35)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":32:32)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":33:18)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":34:18)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":35:32)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:28)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:98)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:64)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:108)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:106)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":36:123)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:36)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:58)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:54)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:31)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:65)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":37:123)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":38:19)
+#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":39:13)
+#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":43:34)
+#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":41:34)
+#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":44:19)
+#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:37)
+#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:42)
+#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:31)
+#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:68)
+#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:60)
+#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":45:119)
+#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":46:36)
+#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":48:20)
+#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":49:20)
+#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":50:35)
+#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:28)
+#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:100)
+#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:65)
+#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:110)
+#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:108)
+#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":51:126)
+#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:37)
+#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:42)
+#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:64)
+#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:60)
+#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:31)
+#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:72)
+#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":52:131)
+#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":53:20)
+#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":56:35)
+#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":55:35)
+#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":57:20)
+#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":59:20)
+#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":60:20)
+#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":61:35)
+#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:28)
+#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:46)
+#loc77 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:38)
+#loc78 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:56)
+#loc79 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:54)
+#loc80 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":62:64)
+#loc81 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:40)
+#loc82 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:36)
+#loc83 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:31)
+#loc84 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:48)
+#loc85 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":63:88)
+#loc86 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":64:20)
+#loc87 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":65:20)
+#loc88 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:25)
+#loc89 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:37)
+#loc90 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/mb/cmb5zlldkrfleyhckhiomhpfsvkfkllotbopalddwkdfhlizveux.py":66:4)
+#loc101 = loc("xoffset"(#loc2))
+#loc102 = loc("xoffset"(#loc3))
+#loc103 = loc("xindex"(#loc4))
+#loc104 = loc("xindex"(#loc5))
+#loc105 = loc("xmask"(#loc6))
+#loc106 = loc("x0"(#loc7))
+#loc107 = loc("x1"(#loc8))
+#loc108 = loc("x1"(#loc9))
+#loc109 = loc("tmp31"(#loc10))
+#loc110 = loc("tmp31"(#loc11))
+#loc111 = loc("tmp31"(#loc12))
+#loc112 = loc("tmp32"(#loc13))
+#loc113 = loc("tmp32"(#loc14))
+#loc114 = loc("tmp1"(#loc15))
+#loc115 = loc("tmp2"(#loc16))
+#loc116 = loc("tmp3"(#loc17))
+#loc117 = loc("tmp3"(#loc18))
+#loc118 = loc("tmp3"(#loc19))
+#loc119 = loc("tmp3"(#loc20))
+#loc120 = loc("tmp3"(#loc21))
+#loc121 = loc("tmp3"(#loc22))
+#loc122 = loc("tmp4"(#loc23))
+#loc123 = loc("tmp5"(#loc24))
+#loc124 = loc("tmp6"(#loc25))
+#loc125 = loc("tmp7"(#loc26))
+#loc126 = loc("tmp8"(#loc27))
+#loc127 = loc("tmp10"(#loc34))
+#loc128 = loc("tmp10"(#loc35))
+#loc129 = loc("tmp10"(#loc36))
+#loc130 = loc("tmp10"(#loc37))
+#loc131 = loc("tmp10"(#loc38))
+#loc132 = loc("tmp10"(#loc39))
+#loc133 = loc("tmp11"(#loc40))
+#loc134 = loc("tmp12"(#loc41))
+#loc135 = loc("tmp16"(#loc42))
+#loc136 = loc("tmp14"(#loc43))
+#loc137 = loc("tmp17"(#loc44))
+#loc138 = loc("tmp18"(#loc45))
+#loc139 = loc("tmp18"(#loc46))
+#loc140 = loc("tmp18"(#loc47))
+#loc141 = loc("tmp18"(#loc48))
+#loc142 = loc("tmp18"(#loc49))
+#loc143 = loc("tmp18"(#loc50))
+#loc144 = loc("tmp19"(#loc51))
+#loc145 = loc("tmp21"(#loc52))
+#loc146 = loc("tmp22"(#loc53))
+#loc147 = loc("tmp23"(#loc54))
+#loc148 = loc("tmp25"(#loc61))
+#loc149 = loc("tmp25"(#loc62))
+#loc150 = loc("tmp25"(#loc63))
+#loc151 = loc("tmp25"(#loc64))
+#loc152 = loc("tmp25"(#loc65))
+#loc153 = loc("tmp25"(#loc66))
+#loc154 = loc("tmp25"(#loc67))
+#loc155 = loc("tmp26"(#loc68))
+#loc156 = loc("tmp29"(#loc69))
+#loc157 = loc("tmp28"(#loc70))
+#loc158 = loc("tmp30"(#loc71))
+#loc159 = loc("tmp34"(#loc72))
+#loc160 = loc("tmp35"(#loc73))
+#loc161 = loc("tmp36"(#loc74))
+#loc162 = loc("tmp38"(#loc81))
+#loc163 = loc("tmp38"(#loc82))
+#loc164 = loc("tmp38"(#loc83))
+#loc165 = loc("tmp38"(#loc84))
+#loc166 = loc("tmp38"(#loc85))
+#loc167 = loc("tmp39"(#loc86))
+#loc168 = loc("tmp40"(#loc87))
+#loc169 = loc(fused[#loc135, #loc136])
+#loc170 = loc(fused[#loc156, #loc157])

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/__grp__triton_red_fused_zeros_0.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_red_fused_zeros_0.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.source", "triton_red_fused_zeros_0.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.ttir", "triton_red_fused_zeros_0.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.ttgir", "triton_red_fused_zeros_0.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.llir", "triton_red_fused_zeros_0.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.ptx", "triton_red_fused_zeros_0.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.cubin", "triton_red_fused_zeros_0.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.json"}}

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.json

@@ -0,0 +1 @@
+{"hash": "34e25b732afeb5d6a31b0c6feb7a22f95462a59f2200623b625166fff2db10d2", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 4, "num_ctas": 1, "num_stages": 1, "warp_size": 32, "maxnreg": null, "cluster_dims": [1, 1, 1], "ptx_version": null, "ptx_options": null, "ir_override": null, "enable_fp_fusion": true, "launch_cooperative_grid": false, "launch_pdl": false, "supported_fp8_dtypes": ["fp8e4b15", "fp8e4nv", "fp8e5"], "deprecated_fp8_dot_operand_dtypes": ["fp8e4b15"], "default_dot_input_precision": "tf32", "allowed_dot_input_precisions": ["tf32", "tf32x3", "ieee"], "max_num_imprecise_acc_default": 1073741824, "extern_libs": [["libdevice", "/workspace/specforge/lib/python3.11/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]], "debug": true, "backend_name": "cuda", "sanitize_overflow": false, "arch": "sm90", "instrumentation_mode": "", "triton_version": "3.5.1", "tensordesc_meta": [], "shared": 256, "tmem_size": 0, "global_scratch_size": 0, "global_scratch_align": 1, "profile_scratch_size": 0, "profile_scratch_align": 1, "name": "triton_red_fused_zeros_0"}

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.llir

@@ -0,0 +1,170 @@
+; ModuleID = 'LLVMDialectModule'
+source_filename = "LLVMDialectModule"
+target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+
+@global_smem = external local_unnamed_addr addrspace(3) global [0 x i8], align 16
+
+; Function Attrs: nounwind
+define ptx_kernel void @triton_red_fused_zeros_0(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, i32 %4, ptr addrspace(1) readnone captures(none) %5, ptr addrspace(1) readnone captures(none) %6) local_unnamed_addr #0 !dbg !4 {
+  %8 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !7
+  %9 = shl i32 %8, 6, !dbg !8
+  %10 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !9
+  %11 = and i32 %10, 126, !dbg !9
+  %12 = lshr exact i32 %11, 1, !dbg !9
+  %13 = or disjoint i32 %12, %9, !dbg !10
+  %14 = shl nuw nsw i32 %10, 2, !dbg !11
+  %15 = and i32 %14, 4, !dbg !11
+  %16 = sdiv i32 %13, 2048, !dbg !12
+  %17 = mul i32 %16, 2048, !dbg !13
+  %.decomposed = sub i32 %13, %17, !dbg !13
+  %18 = srem i32 %16, 32, !dbg !14
+  %19 = sdiv i32 %13, 65536, !dbg !15
+  %20 = shl nsw i32 %18, 7, !dbg !16
+  %21 = shl nsw i32 %.decomposed, 12, !dbg !17
+  %22 = shl i32 %19, 23, !dbg !18
+  %23 = shl i32 %13, 7, !dbg !19
+  %24 = add i32 %22, %21
+  %25 = add i32 %24, %20
+  %26 = zext nneg i32 %15 to i64, !dbg !20
+  %27 = sext i32 %23 to i64, !dbg !20
+  %invariant.gep = getelementptr bfloat, ptr addrspace(1) %1, i64 %27, !dbg !20
+  br label %28, !dbg !20
+
+28:                                               ; preds = %7, %28
+  %indvars.iv = phi i64 [ 0, %7 ], [ %indvars.iv.next, %28 ]
+  %29 = phi float [ 0.000000e+00, %7 ], [ %70, %28 ]
+  %30 = phi float [ 0.000000e+00, %7 ], [ %71, %28 ]
+  %31 = phi float [ 0.000000e+00, %7 ], [ %72, %28 ]
+  %32 = phi float [ 0.000000e+00, %7 ], [ %73, %28 ]
+  %33 = or disjoint i64 %indvars.iv, %26, !dbg !21
+  %34 = trunc nuw nsw i64 %33 to i32, !dbg !22
+  %35 = add i32 %25, %34, !dbg !22
+  %36 = sext i32 %35 to i64, !dbg !23
+  %37 = getelementptr bfloat, ptr addrspace(1) %0, i64 %36, !dbg !23
+  %38 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_first.b64 $0, 1.0;", "=l"() #4, !dbg !24
+  %39 = tail call { i32, i32 } asm sideeffect "mov.u32 $0, $2;\0A\09mov.u32 $1, $3;\0A\09@$6 ld.global.L1::evict_first.L2::cache_hint.v2.b32 { $0, $1 }, [ $4 + 0 ], $5;", "=r,=r,r,r,l,l,b"(i32 0, i32 0, ptr addrspace(1) %37, i64 %38, i1 true) #4, !dbg !24
+  %40 = extractvalue { i32, i32 } %39, 0, !dbg !24
+  %41 = bitcast i32 %40 to <2 x bfloat>, !dbg !24
+  %42 = extractvalue { i32, i32 } %39, 1, !dbg !24
+  %43 = bitcast i32 %42 to <2 x bfloat>, !dbg !24
+  %44 = extractelement <2 x bfloat> %41, i64 0, !dbg !24
+  %45 = extractelement <2 x bfloat> %41, i64 1, !dbg !24
+  %46 = extractelement <2 x bfloat> %43, i64 0, !dbg !24
+  %47 = extractelement <2 x bfloat> %43, i64 1, !dbg !24
+  %48 = fpext bfloat %44 to float, !dbg !25
+  %49 = fpext bfloat %45 to float, !dbg !25
+  %50 = fpext bfloat %46 to float, !dbg !25
+  %51 = fpext bfloat %47 to float, !dbg !25
+  %gep = getelementptr bfloat, ptr addrspace(1) %invariant.gep, i64 %33, !dbg !26
+  %52 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_first.b64 $0, 1.0;", "=l"() #4, !dbg !27
+  %53 = tail call { i32, i32 } asm sideeffect "mov.u32 $0, $2;\0A\09mov.u32 $1, $3;\0A\09@$6 ld.global.L1::evict_first.L2::cache_hint.v2.b32 { $0, $1 }, [ $4 + 0 ], $5;", "=r,=r,r,r,l,l,b"(i32 0, i32 0, ptr addrspace(1) %gep, i64 %52, i1 true) #4, !dbg !27
+  %54 = extractvalue { i32, i32 } %53, 0, !dbg !27
+  %55 = bitcast i32 %54 to <2 x bfloat>, !dbg !27
+  %56 = extractvalue { i32, i32 } %53, 1, !dbg !27
+  %57 = bitcast i32 %56 to <2 x bfloat>, !dbg !27
+  %58 = extractelement <2 x bfloat> %55, i64 0, !dbg !27
+  %59 = extractelement <2 x bfloat> %55, i64 1, !dbg !27
+  %60 = extractelement <2 x bfloat> %57, i64 0, !dbg !27
+  %61 = extractelement <2 x bfloat> %57, i64 1, !dbg !27
+  %62 = fpext bfloat %58 to float, !dbg !28
+  %63 = fpext bfloat %59 to float, !dbg !28
+  %64 = fpext bfloat %60 to float, !dbg !28
+  %65 = fpext bfloat %61 to float, !dbg !28
+  %66 = fmul float %48, %62, !dbg !29
+  %67 = fmul float %49, %63, !dbg !29
+  %68 = fmul float %50, %64, !dbg !29
+  %69 = fmul float %51, %65, !dbg !29
+  %70 = fadd float %29, %66, !dbg !30
+  %71 = fadd float %30, %67, !dbg !30
+  %72 = fadd float %31, %68, !dbg !30
+  %73 = fadd float %32, %69, !dbg !30
+  %indvars.iv.next = add nuw nsw i64 %indvars.iv, 8, !dbg !20
+  %74 = icmp samesign ult i64 %indvars.iv, 120, !dbg !20
+  br i1 %74, label %28, label %75, !dbg !20
+
+75:                                               ; preds = %28
+  %76 = and i32 %10, 63, !dbg !9
+  %77 = or disjoint i32 %9, %76, !dbg !10
+  %78 = fadd float %70, %71, !dbg !31
+  %79 = fadd float %72, %78, !dbg !31
+  %80 = fadd float %73, %79, !dbg !31
+  %81 = bitcast float %80 to i32, !dbg !35
+  %82 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %81, i32 1, i32 31), !dbg !35
+  %83 = bitcast i32 %82 to float, !dbg !35
+  %84 = fadd float %80, %83, !dbg !31
+  %85 = shl nuw nsw i32 %11, 1, !dbg !36
+  %86 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %85, !dbg !36
+  store float %84, ptr addrspace(3) %86, align 4, !dbg !36
+  tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !36
+  %87 = shl nuw nsw i32 %76, 2, !dbg !36
+  %88 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %87, !dbg !36
+  %89 = load i32, ptr addrspace(3) %88, align 4, !dbg !36
+  %90 = sext i32 %77 to i64, !dbg !37
+  %91 = getelementptr float, ptr addrspace(1) %2, i64 %90, !dbg !37
+  %92 = and i32 %10, 64, !dbg !38
+  %93 = icmp eq i32 %92, 0, !dbg !38
+  tail call void asm sideeffect "@$2 st.global.b32 [ $1 + 0 ], { $0 };", "r,l,b"(i32 %89, ptr addrspace(1) %91, i1 %93) #4, !dbg !38
+  ret void, !dbg !39
+}
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
+
+; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
+declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) #2
+
+; Function Attrs: convergent nocallback nounwind
+declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3
+
+attributes #0 = { nounwind "nvvm.reqntid"="128" }
+attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+attributes #2 = { convergent nocallback nounwind memory(inaccessiblemem: readwrite) }
+attributes #3 = { convergent nocallback nounwind }
+attributes #4 = { nounwind }
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
+!1 = !DIFile(filename: "cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og")
+!2 = !{i32 2, !"Debug Info Version", i32 3}
+!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
+!4 = distinct !DISubprogram(name: "triton_red_fused_zeros_0", linkageName: "triton_red_fused_zeros_0", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
+!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
+!6 = !{}
+!7 = !DILocation(line: 23, column: 28, scope: !4)
+!8 = !DILocation(line: 23, column: 33, scope: !4)
+!9 = !DILocation(line: 24, column: 44, scope: !4)
+!10 = !DILocation(line: 24, column: 23, scope: !4)
+!11 = !DILocation(line: 26, column: 37, scope: !4)
+!12 = !DILocation(line: 29, column: 21, scope: !4)
+!13 = !DILocation(line: 28, column: 19, scope: !4)
+!14 = !DILocation(line: 29, column: 29, scope: !4)
+!15 = !DILocation(line: 30, column: 19, scope: !4)
+!16 = !DILocation(line: 39, column: 45, scope: !4)
+!17 = !DILocation(line: 39, column: 55, scope: !4)
+!18 = !DILocation(line: 39, column: 68, scope: !4)
+!19 = !DILocation(line: 40, column: 45, scope: !4)
+!20 = !DILocation(line: 33, column: 40, scope: !4)
+!21 = !DILocation(line: 34, column: 31, scope: !4)
+!22 = !DILocation(line: 39, column: 60, scope: !4)
+!23 = !DILocation(line: 39, column: 34, scope: !4)
+!24 = !DILocation(line: 39, column: 73, scope: !4)
+!25 = !DILocation(line: 39, column: 127, scope: !4)
+!26 = !DILocation(line: 40, column: 34, scope: !4)
+!27 = !DILocation(line: 40, column: 50, scope: !4)
+!28 = !DILocation(line: 40, column: 104, scope: !4)
+!29 = !DILocation(line: 41, column: 22, scope: !4)
+!30 = !DILocation(line: 43, column: 23, scope: !4)
+!31 = !DILocation(line: 261, column: 15, scope: !32, inlinedAt: !34)
+!32 = distinct !DILexicalBlockFile(scope: !4, file: !33, discriminator: 0)
+!33 = !DIFile(filename: "standard.py", directory: "/workspace/specforge/lib/python3.11/site-packages/triton/language")
+!34 = !DILocation(line: 45, column: 25, scope: !4)
+!35 = !DILocation(line: 291, column: 36, scope: !32, inlinedAt: !34)
+!36 = !DILocation(line: 45, column: 28, scope: !4)
+!37 = !DILocation(line: 49, column: 25, scope: !4)
+!38 = !DILocation(line: 49, column: 36, scope: !4)
+!39 = !DILocation(line: 49, column: 4, scope: !4)

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.ptx

@@ -0,0 +1,411 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.7
+.target sm_90a
+.address_size 64
+
+	// .globl	triton_red_fused_zeros_0 // -- Begin function triton_red_fused_zeros_0
+.extern .shared .align 16 .b8 global_smem[];
+                                        // @triton_red_fused_zeros_0
+.visible .entry triton_red_fused_zeros_0(
+	.param .u64 .ptr .global .align 1 triton_red_fused_zeros_0_param_0,
+	.param .u64 .ptr .global .align 1 triton_red_fused_zeros_0_param_1,
+	.param .u64 .ptr .global .align 1 triton_red_fused_zeros_0_param_2,
+	.param .u32 triton_red_fused_zeros_0_param_3,
+	.param .u32 triton_red_fused_zeros_0_param_4,
+	.param .u64 .ptr .global .align 1 triton_red_fused_zeros_0_param_5,
+	.param .u64 .ptr .global .align 1 triton_red_fused_zeros_0_param_6
+)
+.reqntid 128
+{
+	.reg .pred 	%p<5>;
+	.reg .b16 	%rs<9>;
+	.reg .b32 	%r<79>;
+	.reg .b64 	%rd<24>;
+	.loc	1 18 0                          // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:18:0
+$L__func_begin0:
+	.loc	1 18 0                          // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:18:0
+
+// %bb.0:
+	ld.param.b64 	%rd8, [triton_red_fused_zeros_0_param_2];
+	ld.param.b64 	%rd7, [triton_red_fused_zeros_0_param_0];
+	ld.param.b64 	%rd10, [triton_red_fused_zeros_0_param_1];
+$L__tmp0:
+	.loc	1 23 28                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:23:28
+	mov.u32 	%r13, %ctaid.x;
+	.loc	1 23 33                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:23:33
+	shl.b32 	%r1, %r13, 6;
+	.loc	1 24 44                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:24:44
+	mov.u32 	%r2, %tid.x;
+	and.b32 	%r3, %r2, 126;
+	bfe.u32 	%r14, %r2, 1, 6;
+	.loc	1 24 23                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:24:23
+	or.b32 	%r15, %r14, %r1;
+	.loc	1 26 37                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:26:37
+	shl.b32 	%r16, %r2, 2;
+	and.b32 	%r17, %r16, 4;
+	.loc	1 29 21                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:29:21
+	bfe.s32 	%r18, %r13, 25, 1;
+	shr.u32 	%r19, %r18, 21;
+	add.s32 	%r20, %r15, %r19;
+	shr.s32 	%r21, %r20, 11;
+	.loc	1 29 29                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:29:29
+	shr.u32 	%r22, %r21, 27;
+	add.s32 	%r23, %r21, %r22;
+	and.b32 	%r24, %r23, 33554400;
+	sub.s32 	%r25, %r21, %r24;
+	.loc	1 30 19                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:30:19
+	shr.u32 	%r26, %r18, 16;
+	add.s32 	%r27, %r15, %r26;
+	.loc	1 39 45                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:39:45
+	shl.b32 	%r28, %r25, 7;
+	.loc	1 39 68                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:39:68
+	shl.b32 	%r29, %r27, 7;
+	and.b32 	%r30, %r29, -8388608;
+	.loc	1 33 40                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:33:40
+	and.b32 	%r31, %r2, 1;
+	mul.wide.u32 	%rd11, %r31, 8;
+	shl.b32 	%r32, %r13, 13;
+	shl.b32 	%r33, %r14, 7;
+	or.b32 	%r34, %r32, %r33;
+	mul.wide.s32 	%rd12, %r34, 2;
+	or.b64 	%rd13, %rd11, %rd12;
+	add.s64 	%rd22, %rd10, %rd13;
+	shl.b32 	%r35, %r13, 18;
+	add.s32 	%r36, %r30, %r35;
+	shl.b32 	%r37, %r14, 12;
+	or.b32 	%r38, %r36, %r37;
+	add.s32 	%r39, %r38, %r28;
+	or.b32 	%r40, %r39, %r17;
+	shl.b32 	%r41, %r21, 23;
+	sub.s32 	%r42, %r40, %r41;
+	cvt.u64.u32 	%rd2, %r42;
+	mov.b32 	%r75, 0f00000000;
+	mov.b64 	%rd23, -8;
+	mov.b32 	%r76, %r75;
+	mov.b32 	%r77, %r75;
+	mov.b32 	%r78, %r75;
+$L__BB0_1:                              // =>This Inner Loop Header: Depth=1
+	.loc	1 39 34                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:39:34
+	add.s64 	%rd20, %rd2, %rd23;
+	cvt.u32.u64 	%r51, %rd20;
+	add.s32 	%r52, %r51, 8;
+	mad.wide.s32 	%rd15, %r52, 2, %rd7;
+	.loc	1 39 73                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:39:73
+	// begin inline asm
+	mov.u64 %rd14, 0x0;
+	createpolicy.fractional.L2::evict_first.b64 %rd14, 1.0;
+	// end inline asm
+	mov.b32 	%r45, 0;
+	mov.pred 	%p1, -1;
+	// begin inline asm
+	mov.u32 %r43, %r45;
+	mov.u32 %r44, %r45;
+	@%p1 ld.global.L1::evict_first.L2::cache_hint.v2.b32 { %r43, %r44 }, [ %rd15 + 0 ], %rd14;
+	// end inline asm
+	mov.b32 	{%rs1, %rs2}, %r43;
+	mov.b32 	{%rs3, %rs4}, %r44;
+	.loc	1 39 127                        // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:39:127
+	cvt.f32.bf16 	%r53, %rs1;
+	cvt.f32.bf16 	%r54, %rs2;
+	cvt.f32.bf16 	%r55, %rs3;
+	cvt.f32.bf16 	%r56, %rs4;
+	.loc	1 40 50                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:40:50
+	// begin inline asm
+	mov.u64 %rd17, 0x0;
+	createpolicy.fractional.L2::evict_first.b64 %rd17, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u32 %r47, %r45;
+	mov.u32 %r48, %r45;
+	@%p1 ld.global.L1::evict_first.L2::cache_hint.v2.b32 { %r47, %r48 }, [ %rd22 + 0 ], %rd17;
+	// end inline asm
+	mov.b32 	{%rs5, %rs6}, %r47;
+	mov.b32 	{%rs7, %rs8}, %r48;
+	.loc	1 40 104                        // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:40:104
+	cvt.f32.bf16 	%r57, %rs5;
+	cvt.f32.bf16 	%r58, %rs6;
+	cvt.f32.bf16 	%r59, %rs7;
+	cvt.f32.bf16 	%r60, %rs8;
+	.loc	1 43 23                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:43:23
+	fma.rn.f32 	%r75, %r53, %r57, %r75;
+	fma.rn.f32 	%r76, %r54, %r58, %r76;
+	fma.rn.f32 	%r77, %r55, %r59, %r77;
+	fma.rn.f32 	%r78, %r56, %r60, %r78;
+	.loc	1 33 40                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:33:40
+	add.s64 	%rd23, %rd23, 8;
+	add.s64 	%rd22, %rd22, 16;
+	setp.lt.u64 	%p3, %rd23, 120;
+	@%p3 bra 	$L__BB0_1;
+// %bb.2:
+	.loc	1 24 44                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:24:44
+	and.b32 	%r62, %r2, 63;
+	.loc	1 24 23                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:24:23
+	or.b32 	%r63, %r1, %r62;
+$L__tmp1:
+	.loc	2 261 15                        // standard.py:261:15 @[ cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:45:25 ]
+	add.f32 	%r64, %r75, %r76;
+	add.f32 	%r65, %r77, %r64;
+	add.f32 	%r66, %r78, %r65;
+	.loc	2 291 36                        // standard.py:291:36 @[ cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:45:25 ]
+	shfl.sync.bfly.b32 	%r67, %r66, 1, 31, -1;
+	.loc	2 261 15                        // standard.py:261:15 @[ cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:45:25 ]
+	add.f32 	%r68, %r66, %r67;
+$L__tmp2:
+	.loc	1 45 28                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:45:28
+	shl.b32 	%r69, %r3, 1;
+	mov.b32 	%r70, global_smem;
+	add.s32 	%r71, %r70, %r69;
+	st.shared.b32 	[%r71], %r68;
+	bar.sync 	0;
+	shl.b32 	%r72, %r62, 2;
+	add.s32 	%r73, %r70, %r72;
+	ld.shared.b32 	%r61, [%r73];
+	.loc	1 49 25                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:49:25
+	mad.wide.s32 	%rd21, %r63, 4, %rd8;
+	.loc	1 49 36                         // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:49:36
+	and.b32 	%r74, %r2, 64;
+	setp.eq.b32 	%p4, %r74, 0;
+	// begin inline asm
+	@%p4 st.global.b32 [ %rd21 + 0 ], { %r61 };
+	// end inline asm
+	.loc	1 49 4                          // cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py:49:4
+	ret;
+$L__tmp3:
+$L__func_end0:
+                                        // -- End function
+}
+	.file	1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py"
+	.file	2 "/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py"
+	.section	.debug_abbrev
+	{
+.b8 1                                   // Abbreviation Code
+.b8 17                                  // DW_TAG_compile_unit
+.b8 1                                   // DW_CHILDREN_yes
+.b8 37                                  // DW_AT_producer
+.b8 8                                   // DW_FORM_string
+.b8 19                                  // DW_AT_language
+.b8 5                                   // DW_FORM_data2
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 16                                  // DW_AT_stmt_list
+.b8 6                                   // DW_FORM_data4
+.b8 27                                  // DW_AT_comp_dir
+.b8 8                                   // DW_FORM_string
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 2                                   // Abbreviation Code
+.b8 46                                  // DW_TAG_subprogram
+.b8 0                                   // DW_CHILDREN_no
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 32                                  // DW_AT_inline
+.b8 11                                  // DW_FORM_data1
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 3                                   // Abbreviation Code
+.b8 46                                  // DW_TAG_subprogram
+.b8 1                                   // DW_CHILDREN_yes
+.b8 17                                  // DW_AT_low_pc
+.b8 1                                   // DW_FORM_addr
+.b8 18                                  // DW_AT_high_pc
+.b8 1                                   // DW_FORM_addr
+.b8 49                                  // DW_AT_abstract_origin
+.b8 19                                  // DW_FORM_ref4
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 4                                   // Abbreviation Code
+.b8 29                                  // DW_TAG_inlined_subroutine
+.b8 0                                   // DW_CHILDREN_no
+.b8 49                                  // DW_AT_abstract_origin
+.b8 19                                  // DW_FORM_ref4
+.b8 17                                  // DW_AT_low_pc
+.b8 1                                   // DW_FORM_addr
+.b8 18                                  // DW_AT_high_pc
+.b8 1                                   // DW_FORM_addr
+.b8 88                                  // DW_AT_call_file
+.b8 11                                  // DW_FORM_data1
+.b8 89                                  // DW_AT_call_line
+.b8 11                                  // DW_FORM_data1
+.b8 87                                  // DW_AT_call_column
+.b8 11                                  // DW_FORM_data1
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 0                                   // EOM(3)
+	}
+	.section	.debug_info
+	{
+.b32 209                                // Length of Unit
+.b8 2                                   // DWARF version number
+.b8 0
+.b32 .debug_abbrev                      // Offset Into Abbrev. Section
+.b8 8                                   // Address Size (in bytes)
+.b8 1                                   // Abbrev [1] 0xb:0xca DW_TAG_compile_unit
+.b8 116                                 // DW_AT_producer
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2                                   // DW_AT_language
+.b8 0
+.b8 99                                  // DW_AT_name
+.b8 111
+.b8 103
+.b8 111
+.b8 108
+.b8 53
+.b8 53
+.b8 99
+.b8 116
+.b8 104
+.b8 107
+.b8 52
+.b8 122
+.b8 101
+.b8 118
+.b8 115
+.b8 121
+.b8 51
+.b8 100
+.b8 108
+.b8 113
+.b8 105
+.b8 121
+.b8 122
+.b8 105
+.b8 112
+.b8 101
+.b8 102
+.b8 118
+.b8 55
+.b8 51
+.b8 53
+.b8 103
+.b8 101
+.b8 50
+.b8 119
+.b8 116
+.b8 97
+.b8 100
+.b8 100
+.b8 118
+.b8 107
+.b8 52
+.b8 51
+.b8 54
+.b8 113
+.b8 104
+.b8 116
+.b8 53
+.b8 110
+.b8 111
+.b8 120
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line                        // DW_AT_stmt_list
+.b8 47                                  // DW_AT_comp_dir
+.b8 119
+.b8 111
+.b8 114
+.b8 107
+.b8 115
+.b8 112
+.b8 97
+.b8 99
+.b8 101
+.b8 47
+.b8 104
+.b8 97
+.b8 110
+.b8 114
+.b8 117
+.b8 105
+.b8 47
+.b8 83
+.b8 112
+.b8 101
+.b8 99
+.b8 70
+.b8 111
+.b8 114
+.b8 103
+.b8 101
+.b8 45
+.b8 101
+.b8 120
+.b8 116
+.b8 47
+.b8 99
+.b8 97
+.b8 99
+.b8 104
+.b8 101
+.b8 47
+.b8 99
+.b8 111
+.b8 109
+.b8 112
+.b8 105
+.b8 108
+.b8 101
+.b8 100
+.b8 95
+.b8 107
+.b8 101
+.b8 114
+.b8 110
+.b8 101
+.b8 108
+.b8 115
+.b8 47
+.b8 111
+.b8 103
+.b8 0
+.b8 2                                   // Abbrev [2] 0x8b:0x1b DW_TAG_subprogram
+.b8 116                                 // DW_AT_name
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 114
+.b8 101
+.b8 100
+.b8 95
+.b8 102
+.b8 117
+.b8 115
+.b8 101
+.b8 100
+.b8 95
+.b8 122
+.b8 101
+.b8 114
+.b8 111
+.b8 115
+.b8 95
+.b8 48
+.b8 0
+.b8 1                                   // DW_AT_inline
+.b8 3                                   // Abbrev [3] 0xa6:0x2e DW_TAG_subprogram
+.b64 $L__func_begin0                    // DW_AT_low_pc
+.b64 $L__func_end0                      // DW_AT_high_pc
+.b32 139                                // DW_AT_abstract_origin
+.b8 4                                   // Abbrev [4] 0xbb:0x18 DW_TAG_inlined_subroutine
+.b32 139                                // DW_AT_abstract_origin
+.b64 $L__tmp1                           // DW_AT_low_pc
+.b64 $L__tmp2                           // DW_AT_high_pc
+.b8 1                                   // DW_AT_call_file
+.b8 45                                  // DW_AT_call_line
+.b8 25                                  // DW_AT_call_column
+.b8 0                                   // End Of Children Mark
+.b8 0                                   // End Of Children Mark
+	}
+	.section	.debug_macinfo	{	}

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.source

@@ -0,0 +1,222 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":18:0)
+#loc44 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":285:0)
+#loc46 = loc(unknown)
+#loc49 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":260:0)
+#loc53 = loc("in_ptr0"(#loc))
+#loc54 = loc("in_ptr1"(#loc))
+#loc55 = loc("out_ptr1"(#loc))
+#loc56 = loc("xnumel"(#loc))
+#loc57 = loc("r0_numel"(#loc))
+#loc97 = loc("input"(#loc44))
+#loc98 = loc("a"(#loc49))
+#loc99 = loc("b"(#loc49))
+module {
+  tt.func public @triton_red_fused_zeros_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %out_ptr1: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr1"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %xnumel_0 = arith.constant 131072 : i32 loc(#loc58)
+    %r0_numel_1 = arith.constant 128 : i32 loc(#loc59)
+    %xoffset = tt.get_program_id x : i32 loc(#loc60)
+    %xoffset_2 = arith.constant 64 : i32 loc(#loc61)
+    %xoffset_3 = arith.constant 64 : i32 loc(#loc61)
+    %xoffset_4 = arith.muli %xoffset, %xoffset_3 : i32 loc(#loc61)
+    %xindex = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32> loc(#loc62)
+    %xindex_5 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<64xi32> -> tensor<64x1xi32> loc(#loc63)
+    %xindex_6 = tt.splat %xoffset_4 : i32 -> tensor<64x1xi32> loc(#loc64)
+    %xindex_7 = arith.addi %xindex_6, %xindex_5 : tensor<64x1xi32> loc(#loc64)
+    %xmask = arith.constant true loc(#loc65)
+    %xmask_8 = arith.constant dense<true> : tensor<64x8xi1> loc(#loc65)
+    %r0_base = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc66)
+    %r0_base_9 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<8xi32> -> tensor<1x8xi32> loc(#loc67)
+    %x0 = arith.constant 2048 : i32 loc(#loc68)
+    %x0_10 = arith.constant 2048 : i32 loc(#loc68)
+    %x0_11 = arith.constant dense<2048> : tensor<64x1xi32> loc(#loc68)
+    %x0_12 = arith.remsi %xindex_7, %x0_11 : tensor<64x1xi32> loc(#loc68)
+    %x1 = arith.constant 2048 : i32 loc(#loc69)
+    %x1_13 = arith.constant 2048 : i32 loc(#loc69)
+    %x1_14 = arith.constant dense<2048> : tensor<64x1xi32> loc(#loc69)
+    %x1_15 = arith.divsi %xindex_7, %x1_14 : tensor<64x1xi32> loc(#loc69)
+    %x1_16 = arith.constant 32 : i32 loc(#loc70)
+    %x1_17 = arith.constant 32 : i32 loc(#loc70)
+    %x1_18 = arith.constant dense<32> : tensor<64x1xi32> loc(#loc70)
+    %x1_19 = arith.remsi %x1_15, %x1_18 : tensor<64x1xi32> loc(#loc70)
+    %x2 = arith.constant 65536 : i32 loc(#loc71)
+    %x2_20 = arith.constant 65536 : i32 loc(#loc71)
+    %x2_21 = arith.constant dense<65536> : tensor<64x1xi32> loc(#loc71)
+    %x2_22 = arith.divsi %xindex_7, %x2_21 : tensor<64x1xi32> loc(#loc71)
+    %_tmp4 = arith.constant 0.000000e+00 : f32 loc(#loc72)
+    %_tmp4_23 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc72)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc16)
+    %c8_i32 = arith.constant 8 : i32 loc(#loc16)
+    %0 = arith.bitcast %c0_i32 : i32 to i32 loc(#loc16)
+    %1 = arith.bitcast %r0_numel_1 : i32 to i32 loc(#loc16)
+    %2 = arith.bitcast %c8_i32 : i32 to i32 loc(#loc16)
+    %3 = ub.poison : i32 loc(#loc16)
+    %_tmp4_24 = scf.for %r0_offset = %0 to %1 step %2 iter_args(%_tmp4_27 = %_tmp4_23) -> (tensor<64x8xf32>)  : i32 {
+      %r0_index = tt.splat %r0_offset : i32 -> tensor<1x8xi32> loc(#loc74)
+      %r0_index_28 = arith.addi %r0_index, %r0_base_9 : tensor<1x8xi32> loc(#loc74)
+      %r0_mask = arith.constant dense<128> : tensor<1x8xi32> loc(#loc75)
+      %r0_mask_29 = arith.cmpi slt, %r0_index_28, %r0_mask : tensor<1x8xi32> loc(#loc75)
+      %tmp0 = arith.constant 128 : i32 loc(#loc76)
+      %tmp0_30 = arith.constant 128 : i32 loc(#loc76)
+      %tmp0_31 = arith.constant dense<128> : tensor<64x1xi32> loc(#loc76)
+      %tmp0_32 = arith.muli %tmp0_31, %x1_19 : tensor<64x1xi32> loc(#loc76)
+      %tmp0_33 = tt.broadcast %r0_index_28 : tensor<1x8xi32> -> tensor<64x8xi32> loc(#loc77)
+      %tmp0_34 = tt.broadcast %tmp0_32 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc77)
+      %tmp0_35 = arith.addi %tmp0_33, %tmp0_34 : tensor<64x8xi32> loc(#loc77)
+      %tmp0_36 = arith.constant 4096 : i32 loc(#loc78)
+      %tmp0_37 = arith.constant 4096 : i32 loc(#loc78)
+      %tmp0_38 = arith.constant dense<4096> : tensor<64x1xi32> loc(#loc78)
+      %tmp0_39 = arith.muli %tmp0_38, %x0_12 : tensor<64x1xi32> loc(#loc78)
+      %tmp0_40 = tt.broadcast %tmp0_39 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc79)
+      %tmp0_41 = arith.addi %tmp0_35, %tmp0_40 : tensor<64x8xi32> loc(#loc79)
+      %tmp0_42 = arith.constant 8388608 : i32 loc(#loc80)
+      %tmp0_43 = arith.constant 8388608 : i32 loc(#loc80)
+      %tmp0_44 = arith.constant dense<8388608> : tensor<64x1xi32> loc(#loc80)
+      %tmp0_45 = arith.muli %tmp0_44, %x2_22 : tensor<64x1xi32> loc(#loc80)
+      %tmp0_46 = tt.broadcast %tmp0_45 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc81)
+      %tmp0_47 = arith.addi %tmp0_41, %tmp0_46 : tensor<64x8xi32> loc(#loc81)
+      %tmp0_48 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc82)
+      %tmp0_49 = tt.addptr %tmp0_48, %tmp0_47 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi32> loc(#loc82)
+      %tmp0_50 = arith.constant 0.000000e+00 : f32 loc(#loc83)
+      %tmp0_51 = tt.broadcast %r0_mask_29 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc83)
+      %tmp0_52 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc83)
+      %tmp0_53 = arith.truncf %tmp0_52 : tensor<64x8xf32> to tensor<64x8xbf16> loc(#loc83)
+      %tmp0_54 = tt.load %tmp0_49, %tmp0_51, %tmp0_53 evictionPolicy = evict_first : tensor<64x8x!tt.ptr<bf16>> loc(#loc83)
+      %tmp0_55 = arith.extf %tmp0_54 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc84)
+      %tmp1 = arith.constant 128 : i32 loc(#loc85)
+      %tmp1_56 = arith.constant 128 : i32 loc(#loc85)
+      %tmp1_57 = arith.constant dense<128> : tensor<64x1xi32> loc(#loc85)
+      %tmp1_58 = arith.muli %tmp1_57, %xindex_7 : tensor<64x1xi32> loc(#loc85)
+      %tmp1_59 = tt.broadcast %r0_index_28 : tensor<1x8xi32> -> tensor<64x8xi32> loc(#loc86)
+      %tmp1_60 = tt.broadcast %tmp1_58 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc86)
+      %tmp1_61 = arith.addi %tmp1_59, %tmp1_60 : tensor<64x8xi32> loc(#loc86)
+      %tmp1_62 = tt.splat %in_ptr1 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc87)
+      %tmp1_63 = tt.addptr %tmp1_62, %tmp1_61 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi32> loc(#loc87)
+      %tmp1_64 = arith.constant 0.000000e+00 : f32 loc(#loc88)
+      %tmp1_65 = tt.broadcast %r0_mask_29 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc88)
+      %tmp1_66 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc88)
+      %tmp1_67 = arith.truncf %tmp1_66 : tensor<64x8xf32> to tensor<64x8xbf16> loc(#loc88)
+      %tmp1_68 = tt.load %tmp1_63, %tmp1_65, %tmp1_67 evictionPolicy = evict_first : tensor<64x8x!tt.ptr<bf16>> loc(#loc88)
+      %tmp1_69 = arith.extf %tmp1_68 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc89)
+      %tmp2 = arith.mulf %tmp0_55, %tmp1_69 : tensor<64x8xf32> loc(#loc90)
+      %tmp5 = arith.addf %_tmp4_27, %tmp2 : tensor<64x8xf32> loc(#loc91)
+      %_tmp4_70 = tt.broadcast %r0_mask_29 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc92)
+      %_tmp4_71 = arith.select %_tmp4_70, %tmp5, %_tmp4_27 : tensor<64x8xi1>, tensor<64x8xf32> loc(#loc92)
+      scf.yield %_tmp4_71 : tensor<64x8xf32> loc(#loc36)
+    } loc(#loc73)
+    %tmp4 = tt.call @"triton.language.standard.sum__fp32S64_8S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%_tmp4_24) : (tensor<64x8xf32>) -> tensor<64xf32> loc(#loc93)
+    %tmp4_25 = tt.expand_dims %tmp4 {axis = 1 : i32} : tensor<64xf32> -> tensor<64x1xf32> loc(#loc94)
+    %tmp7 = arith.constant 0.000000e+00 : f32 loc(#loc95)
+    %tmp8 = arith.constant dense<0.000000e+00> : tensor<64x1xf32> loc(#loc96)
+    %tmp8_26 = arith.subf %tmp4_25, %tmp8 : tensor<64x1xf32> loc(#loc96)
+    %4 = tt.splat %out_ptr1 : !tt.ptr<f32> -> tensor<64x1x!tt.ptr<f32>> loc(#loc41)
+    %5 = tt.addptr %4, %xindex_7 : tensor<64x1x!tt.ptr<f32>>, tensor<64x1xi32> loc(#loc41)
+    tt.store %5, %tmp8_26 : tensor<64x1x!tt.ptr<f32>> loc(#loc42)
+    tt.return loc(#loc43)
+  } loc(#loc)
+  tt.func private @"triton.language.standard.sum__fp32S64_8S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%input: tensor<64x8xf32> loc("input"(#loc44))) -> tensor<64xf32> attributes {noinline = false} {
+    %0 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
+    ^bb0(%arg1: f32 loc(unknown), %arg2: f32 loc(unknown)):
+      %2 = tt.call @triton.language.standard._sum_combine__fp32_fp32__(%arg1, %arg2) : (f32, f32) -> f32 loc(#loc45)
+      tt.reduce.return %2 : f32 loc(#loc45)
+    }) : (tensor<64x8xf32>) -> tensor<64xf32> loc(#loc45)
+    tt.return %0 : tensor<64xf32> loc(#loc47)
+  ^bb1:  // no predecessors
+    %1 = ub.poison : tensor<64xf32> loc(#loc48)
+    tt.return %1 : tensor<64xf32> loc(#loc48)
+  } loc(#loc44)
+  tt.func private @triton.language.standard._sum_combine__fp32_fp32__(%a: f32 loc("a"(#loc49)), %b: f32 loc("b"(#loc49))) -> f32 attributes {noinline = false} {
+    %0 = arith.addf %a, %b : f32 loc(#loc50)
+    tt.return %0 : f32 loc(#loc51)
+  ^bb1:  // no predecessors
+    %1 = ub.poison : f32 loc(#loc52)
+    tt.return %1 : f32 loc(#loc52)
+  } loc(#loc49)
+} loc(#loc)
+#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":19:13)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":20:15)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":23:28)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":23:33)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:36)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:44)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:23)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":25:46)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":26:27)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":26:37)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":28:19)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":29:21)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":29:29)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":30:19)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":32:43)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":33:40)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":34:31)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":35:29)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:45)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:41)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:55)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:50)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:68)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:60)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:34)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:73)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:127)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:45)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:41)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:34)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:50)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:104)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":41:22)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":43:23)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":44:40)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":44:8)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":45:25)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":45:28)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":47:11)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":48:18)
+#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:25)
+#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:36)
+#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:4)
+#loc45 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc47 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:11)
+#loc48 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:4)
+#loc50 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc51 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:11)
+#loc52 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:4)
+#loc58 = loc("xnumel"(#loc1))
+#loc59 = loc("r0_numel"(#loc2))
+#loc60 = loc("xoffset"(#loc3))
+#loc61 = loc("xoffset"(#loc4))
+#loc62 = loc("xindex"(#loc5))
+#loc63 = loc("xindex"(#loc6))
+#loc64 = loc("xindex"(#loc7))
+#loc65 = loc("xmask"(#loc8))
+#loc66 = loc("r0_base"(#loc9))
+#loc67 = loc("r0_base"(#loc10))
+#loc68 = loc("x0"(#loc11))
+#loc69 = loc("x1"(#loc12))
+#loc70 = loc("x1"(#loc13))
+#loc71 = loc("x2"(#loc14))
+#loc72 = loc("_tmp4"(#loc15))
+#loc73 = loc("_tmp4"(#loc16))
+#loc74 = loc("r0_index"(#loc17))
+#loc75 = loc("r0_mask"(#loc18))
+#loc76 = loc("tmp0"(#loc19))
+#loc77 = loc("tmp0"(#loc20))
+#loc78 = loc("tmp0"(#loc21))
+#loc79 = loc("tmp0"(#loc22))
+#loc80 = loc("tmp0"(#loc23))
+#loc81 = loc("tmp0"(#loc24))
+#loc82 = loc("tmp0"(#loc25))
+#loc83 = loc("tmp0"(#loc26))
+#loc84 = loc("tmp0"(#loc27))
+#loc85 = loc("tmp1"(#loc28))
+#loc86 = loc("tmp1"(#loc29))
+#loc87 = loc("tmp1"(#loc30))
+#loc88 = loc("tmp1"(#loc31))
+#loc89 = loc("tmp1"(#loc32))
+#loc90 = loc("tmp2"(#loc33))
+#loc91 = loc("tmp5"(#loc34))
+#loc92 = loc("_tmp4"(#loc35))
+#loc93 = loc("tmp4"(#loc37))
+#loc94 = loc("tmp4"(#loc38))
+#loc95 = loc("tmp7"(#loc39))
+#loc96 = loc("tmp8"(#loc40))

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.ttgir

@@ -0,0 +1,155 @@
+#blocked = #ttg.blocked<{sizePerThread = [1, 4], threadsPerWarp = [16, 2], warpsPerCTA = [4, 1], order = [1, 0]}>
+#blocked1 = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [2, 2], order = [0, 1]}>
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":18:0)
+#loc1 = loc(unknown)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":45:25)
+#loc39 = loc("in_ptr0"(#loc))
+#loc40 = loc("in_ptr1"(#loc))
+#loc41 = loc("out_ptr1"(#loc))
+#loc42 = loc("xnumel"(#loc))
+#loc43 = loc("r0_numel"(#loc))
+#loc73 = loc("tmp4"(#loc33))
+#loc76 = loc(callsite(#loc1 at #loc73))
+module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
+  tt.func public @triton_red_fused_zeros_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %out_ptr1: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr1"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<128> : tensor<1x8xi32, #blocked> loc(#loc1)
+    %cst_0 = arith.constant dense<128> : tensor<64x1xi32, #blocked> loc(#loc1)
+    %cst_1 = arith.constant dense<4096> : tensor<64x1xi32, #blocked> loc(#loc1)
+    %cst_2 = arith.constant dense<8388608> : tensor<64x1xi32, #blocked> loc(#loc1)
+    %cst_3 = arith.constant dense<65536> : tensor<64x1xi32, #blocked> loc(#loc1)
+    %cst_4 = arith.constant dense<32> : tensor<64x1xi32, #blocked> loc(#loc1)
+    %cst_5 = arith.constant dense<2048> : tensor<64x1xi32, #blocked> loc(#loc1)
+    %c64_i32 = arith.constant 64 : i32 loc(#loc1)
+    %cst_6 = arith.constant dense<0.000000e+00> : tensor<64x8xbf16, #blocked> loc(#loc1)
+    %c8_i32 = arith.constant 8 : i32 loc(#loc1)
+    %c128_i32 = arith.constant 128 : i32 loc(#loc1)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc1)
+    %cst_7 = arith.constant dense<0.000000e+00> : tensor<64x8xf32, #blocked> loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc44)
+    %xoffset_8 = arith.muli %xoffset, %c64_i32 : i32 loc(#loc45)
+    %xindex = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc46)
+    %xindex_9 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> loc(#loc46)
+    %xindex_10 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<64x1xi32, #blocked> loc(#loc46)
+    %xindex_11 = tt.expand_dims %xindex_9 {axis = 1 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> -> tensor<64x1xi32, #blocked1> loc(#loc46)
+    %xindex_12 = tt.splat %xoffset_8 : i32 -> tensor<64x1xi32, #blocked> loc(#loc47)
+    %xindex_13 = tt.splat %xoffset_8 : i32 -> tensor<64x1xi32, #blocked1> loc(#loc47)
+    %xindex_14 = arith.addi %xindex_12, %xindex_10 : tensor<64x1xi32, #blocked> loc(#loc47)
+    %xindex_15 = arith.addi %xindex_13, %xindex_11 : tensor<64x1xi32, #blocked1> loc(#loc47)
+    %r0_base = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc48)
+    %r0_base_16 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<8xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x8xi32, #blocked> loc(#loc48)
+    %x0 = arith.remsi %xindex_14, %cst_5 : tensor<64x1xi32, #blocked> loc(#loc49)
+    %x1 = arith.divsi %xindex_14, %cst_5 : tensor<64x1xi32, #blocked> loc(#loc50)
+    %x1_17 = arith.remsi %x1, %cst_4 : tensor<64x1xi32, #blocked> loc(#loc51)
+    %x2 = arith.divsi %xindex_14, %cst_3 : tensor<64x1xi32, #blocked> loc(#loc52)
+    %tmp0 = arith.muli %x1_17, %cst_0 : tensor<64x1xi32, #blocked> loc(#loc53)
+    %tmp0_18 = tt.broadcast %tmp0 : tensor<64x1xi32, #blocked> -> tensor<64x8xi32, #blocked> loc(#loc54)
+    %tmp0_19 = arith.muli %x0, %cst_1 : tensor<64x1xi32, #blocked> loc(#loc55)
+    %tmp0_20 = tt.broadcast %tmp0_19 : tensor<64x1xi32, #blocked> -> tensor<64x8xi32, #blocked> loc(#loc56)
+    %tmp0_21 = arith.muli %x2, %cst_2 : tensor<64x1xi32, #blocked> loc(#loc57)
+    %tmp0_22 = tt.broadcast %tmp0_21 : tensor<64x1xi32, #blocked> -> tensor<64x8xi32, #blocked> loc(#loc58)
+    %tmp0_23 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc59)
+    %tmp1 = arith.muli %xindex_14, %cst_0 : tensor<64x1xi32, #blocked> loc(#loc60)
+    %tmp1_24 = tt.broadcast %tmp1 : tensor<64x1xi32, #blocked> -> tensor<64x8xi32, #blocked> loc(#loc61)
+    %tmp1_25 = tt.splat %in_ptr1 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc62)
+    %_tmp4 = scf.for %_tmp4_28 = %c0_i32 to %c128_i32 step %c8_i32 iter_args(%arg6 = %cst_7) -> (tensor<64x8xf32, #blocked>)  : i32 {
+      %r0_index = tt.splat %_tmp4_28 : i32 -> tensor<1x8xi32, #blocked> loc(#loc64)
+      %r0_index_29 = arith.addi %r0_index, %r0_base_16 : tensor<1x8xi32, #blocked> loc(#loc64)
+      %r0_mask = arith.cmpi slt, %r0_index_29, %cst : tensor<1x8xi32, #blocked> loc(#loc65)
+      %tmp0_30 = tt.broadcast %r0_index_29 : tensor<1x8xi32, #blocked> -> tensor<64x8xi32, #blocked> loc(#loc54)
+      %tmp0_31 = arith.addi %tmp0_30, %tmp0_18 : tensor<64x8xi32, #blocked> loc(#loc54)
+      %tmp0_32 = arith.addi %tmp0_31, %tmp0_20 : tensor<64x8xi32, #blocked> loc(#loc56)
+      %tmp0_33 = arith.addi %tmp0_32, %tmp0_22 : tensor<64x8xi32, #blocked> loc(#loc58)
+      %tmp0_34 = tt.addptr %tmp0_23, %tmp0_33 : tensor<64x8x!tt.ptr<bf16>, #blocked>, tensor<64x8xi32, #blocked> loc(#loc59)
+      %tmp0_35 = tt.broadcast %r0_mask : tensor<1x8xi1, #blocked> -> tensor<64x8xi1, #blocked> loc(#loc66)
+      %tmp0_36 = tt.load %tmp0_34, %tmp0_35, %cst_6 evictionPolicy = evict_first : tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc66)
+      %tmp0_37 = arith.extf %tmp0_36 : tensor<64x8xbf16, #blocked> to tensor<64x8xf32, #blocked> loc(#loc67)
+      %tmp1_38 = arith.addi %tmp0_30, %tmp1_24 : tensor<64x8xi32, #blocked> loc(#loc61)
+      %tmp1_39 = tt.addptr %tmp1_25, %tmp1_38 : tensor<64x8x!tt.ptr<bf16>, #blocked>, tensor<64x8xi32, #blocked> loc(#loc62)
+      %tmp1_40 = tt.load %tmp1_39, %tmp0_35, %cst_6 evictionPolicy = evict_first : tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc68)
+      %tmp1_41 = arith.extf %tmp1_40 : tensor<64x8xbf16, #blocked> to tensor<64x8xf32, #blocked> loc(#loc69)
+      %tmp2 = arith.mulf %tmp0_37, %tmp1_41 : tensor<64x8xf32, #blocked> loc(#loc70)
+      %tmp5 = arith.addf %arg6, %tmp2 : tensor<64x8xf32, #blocked> loc(#loc71)
+      %_tmp4_42 = arith.select %tmp0_35, %tmp5, %arg6 : tensor<64x8xi1, #blocked>, tensor<64x8xf32, #blocked> loc(#loc72)
+      scf.yield %_tmp4_42 : tensor<64x8xf32, #blocked> loc(#loc31)
+    } loc(#loc63)
+    %tmp4 = "tt.reduce"(%_tmp4) <{axis = 1 : i32}> ({
+    ^bb0(%tmp4_28: f32 loc(callsite(#loc1 at #loc73)), %tmp4_29: f32 loc(callsite(#loc1 at #loc73))):
+      %tmp4_30 = arith.addf %tmp4_28, %tmp4_29 : f32 loc(#loc77)
+      tt.reduce.return %tmp4_30 : f32 loc(#loc75)
+    }) : (tensor<64x8xf32, #blocked>) -> tensor<64xf32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc75)
+    %tmp4_26 = ttg.convert_layout %tmp4 : tensor<64xf32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<64xf32, #ttg.slice<{dim = 1, parent = #blocked1}>> loc(#loc74)
+    %tmp4_27 = tt.expand_dims %tmp4_26 {axis = 1 : i32} : tensor<64xf32, #ttg.slice<{dim = 1, parent = #blocked1}>> -> tensor<64x1xf32, #blocked1> loc(#loc74)
+    %0 = tt.splat %out_ptr1 : !tt.ptr<f32> -> tensor<64x1x!tt.ptr<f32>, #blocked1> loc(#loc36)
+    %1 = tt.addptr %0, %xindex_15 : tensor<64x1x!tt.ptr<f32>, #blocked1>, tensor<64x1xi32, #blocked1> loc(#loc36)
+    tt.store %1, %tmp4_27 : tensor<64x1x!tt.ptr<f32>, #blocked1> loc(#loc37)
+    tt.return loc(#loc38)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":23:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":23:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:44)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":26:37)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":28:19)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":29:21)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":29:29)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":30:19)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:45)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:41)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:55)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:50)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:68)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:60)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:34)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:45)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:41)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:34)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":33:40)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":34:31)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":35:29)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:73)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:127)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:50)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:104)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":41:22)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":43:23)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":44:40)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":44:8)
+#loc32 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":45:28)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:25)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:36)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:4)
+#loc44 = loc("xoffset"(#loc2))
+#loc45 = loc("xoffset"(#loc3))
+#loc46 = loc("xindex"(#loc4))
+#loc47 = loc("xindex"(#loc5))
+#loc48 = loc("r0_base"(#loc6))
+#loc49 = loc("x0"(#loc7))
+#loc50 = loc("x1"(#loc8))
+#loc51 = loc("x1"(#loc9))
+#loc52 = loc("x2"(#loc10))
+#loc53 = loc("tmp0"(#loc11))
+#loc54 = loc("tmp0"(#loc12))
+#loc55 = loc("tmp0"(#loc13))
+#loc56 = loc("tmp0"(#loc14))
+#loc57 = loc("tmp0"(#loc15))
+#loc58 = loc("tmp0"(#loc16))
+#loc59 = loc("tmp0"(#loc17))
+#loc60 = loc("tmp1"(#loc18))
+#loc61 = loc("tmp1"(#loc19))
+#loc62 = loc("tmp1"(#loc20))
+#loc63 = loc("_tmp4"(#loc21))
+#loc64 = loc("r0_index"(#loc22))
+#loc65 = loc("r0_mask"(#loc23))
+#loc66 = loc("tmp0"(#loc24))
+#loc67 = loc("tmp0"(#loc25))
+#loc68 = loc("tmp1"(#loc26))
+#loc69 = loc("tmp1"(#loc27))
+#loc70 = loc("tmp2"(#loc28))
+#loc71 = loc("tmp5"(#loc29))
+#loc72 = loc("_tmp4"(#loc30))
+#loc74 = loc("tmp4"(#loc35))
+#loc75 = loc(callsite(#loc32 at #loc73))
+#loc77 = loc(callsite(#loc34 at #loc75))

SpecForge-ext/cache/compiled_kernels/triton/6/GTRFW4ZK7225NIY3BRX6W6RC7FKGFJM7EIAGEO3CKFTP74W3CDJA/triton_red_fused_zeros_0.ttir

@@ -0,0 +1,152 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":18:0)
+#loc1 = loc(unknown)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":45:25)
+#loc41 = loc("in_ptr0"(#loc))
+#loc42 = loc("in_ptr1"(#loc))
+#loc43 = loc("out_ptr1"(#loc))
+#loc44 = loc("xnumel"(#loc))
+#loc45 = loc("r0_numel"(#loc))
+#loc77 = loc("tmp4"(#loc35))
+#loc80 = loc(callsite(#loc1 at #loc77))
+module {
+  tt.func public @triton_red_fused_zeros_0(%in_ptr0: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<bf16> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %out_ptr1: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr1"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 {tt.divisibility = 16 : i32} loc("r0_numel"(#loc))) attributes {noinline = false} {
+    %cst = arith.constant dense<0.000000e+00> : tensor<64x8xbf16> loc(#loc1)
+    %c8_i32 = arith.constant 8 : i32 loc(#loc2)
+    %c128_i32 = arith.constant 128 : i32 loc(#loc2)
+    %c0_i32 = arith.constant 0 : i32 loc(#loc2)
+    %cst_0 = arith.constant dense<8388608> : tensor<64x1xi32> loc(#loc1)
+    %cst_1 = arith.constant dense<4096> : tensor<64x1xi32> loc(#loc1)
+    %cst_2 = arith.constant dense<128> : tensor<64x1xi32> loc(#loc1)
+    %cst_3 = arith.constant dense<128> : tensor<1x8xi32> loc(#loc1)
+    %cst_4 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc1)
+    %x2 = arith.constant dense<65536> : tensor<64x1xi32> loc(#loc46)
+    %x1 = arith.constant dense<32> : tensor<64x1xi32> loc(#loc47)
+    %cst_5 = arith.constant dense<2048> : tensor<64x1xi32> loc(#loc1)
+    %c64_i32 = arith.constant 64 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc48)
+    %xoffset_6 = arith.muli %xoffset, %c64_i32 : i32 loc(#loc49)
+    %xindex = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32> loc(#loc50)
+    %xindex_7 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<64xi32> -> tensor<64x1xi32> loc(#loc51)
+    %xindex_8 = tt.splat %xoffset_6 : i32 -> tensor<64x1xi32> loc(#loc52)
+    %xindex_9 = arith.addi %xindex_8, %xindex_7 : tensor<64x1xi32> loc(#loc52)
+    %r0_base = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc53)
+    %r0_base_10 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<8xi32> -> tensor<1x8xi32> loc(#loc54)
+    %x0 = arith.remsi %xindex_9, %cst_5 : tensor<64x1xi32> loc(#loc55)
+    %x1_11 = arith.divsi %xindex_9, %cst_5 : tensor<64x1xi32> loc(#loc56)
+    %x1_12 = arith.remsi %x1_11, %x1 : tensor<64x1xi32> loc(#loc47)
+    %x2_13 = arith.divsi %xindex_9, %x2 : tensor<64x1xi32> loc(#loc46)
+    %_tmp4 = scf.for %r0_offset = %c0_i32 to %c128_i32 step %c8_i32 iter_args(%_tmp4_15 = %cst_4) -> (tensor<64x8xf32>)  : i32 {
+      %r0_index = tt.splat %r0_offset : i32 -> tensor<1x8xi32> loc(#loc58)
+      %r0_index_16 = arith.addi %r0_index, %r0_base_10 : tensor<1x8xi32> loc(#loc58)
+      %r0_mask = arith.cmpi slt, %r0_index_16, %cst_3 : tensor<1x8xi32> loc(#loc59)
+      %tmp0 = arith.muli %x1_12, %cst_2 : tensor<64x1xi32> loc(#loc60)
+      %tmp0_17 = tt.broadcast %r0_index_16 : tensor<1x8xi32> -> tensor<64x8xi32> loc(#loc61)
+      %tmp0_18 = tt.broadcast %tmp0 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc61)
+      %tmp0_19 = arith.addi %tmp0_17, %tmp0_18 : tensor<64x8xi32> loc(#loc61)
+      %tmp0_20 = arith.muli %x0, %cst_1 : tensor<64x1xi32> loc(#loc62)
+      %tmp0_21 = tt.broadcast %tmp0_20 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc63)
+      %tmp0_22 = arith.addi %tmp0_19, %tmp0_21 : tensor<64x8xi32> loc(#loc63)
+      %tmp0_23 = arith.muli %x2_13, %cst_0 : tensor<64x1xi32> loc(#loc64)
+      %tmp0_24 = tt.broadcast %tmp0_23 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc65)
+      %tmp0_25 = arith.addi %tmp0_22, %tmp0_24 : tensor<64x8xi32> loc(#loc65)
+      %tmp0_26 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc66)
+      %tmp0_27 = tt.addptr %tmp0_26, %tmp0_25 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi32> loc(#loc66)
+      %tmp0_28 = tt.broadcast %r0_mask : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc67)
+      %tmp0_29 = tt.load %tmp0_27, %tmp0_28, %cst evictionPolicy = evict_first : tensor<64x8x!tt.ptr<bf16>> loc(#loc67)
+      %tmp0_30 = arith.extf %tmp0_29 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc68)
+      %tmp1 = arith.muli %xindex_9, %cst_2 : tensor<64x1xi32> loc(#loc69)
+      %tmp1_31 = tt.broadcast %tmp1 : tensor<64x1xi32> -> tensor<64x8xi32> loc(#loc70)
+      %tmp1_32 = arith.addi %tmp0_17, %tmp1_31 : tensor<64x8xi32> loc(#loc70)
+      %tmp1_33 = tt.splat %in_ptr1 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc71)
+      %tmp1_34 = tt.addptr %tmp1_33, %tmp1_32 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi32> loc(#loc71)
+      %tmp1_35 = tt.load %tmp1_34, %tmp0_28, %cst evictionPolicy = evict_first : tensor<64x8x!tt.ptr<bf16>> loc(#loc72)
+      %tmp1_36 = arith.extf %tmp1_35 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc73)
+      %tmp2 = arith.mulf %tmp0_30, %tmp1_36 : tensor<64x8xf32> loc(#loc74)
+      %tmp5 = arith.addf %_tmp4_15, %tmp2 : tensor<64x8xf32> loc(#loc75)
+      %_tmp4_37 = arith.select %tmp0_28, %tmp5, %_tmp4_15 : tensor<64x8xi1>, tensor<64x8xf32> loc(#loc76)
+      scf.yield %_tmp4_37 : tensor<64x8xf32> loc(#loc33)
+    } loc(#loc57)
+    %tmp4 = "tt.reduce"(%_tmp4) <{axis = 1 : i32}> ({
+    ^bb0(%tmp4_15: f32 loc(callsite(#loc1 at #loc77)), %tmp4_16: f32 loc(callsite(#loc1 at #loc77))):
+      %tmp4_17 = arith.addf %tmp4_15, %tmp4_16 : f32 loc(#loc81)
+      tt.reduce.return %tmp4_17 : f32 loc(#loc79)
+    }) : (tensor<64x8xf32>) -> tensor<64xf32> loc(#loc79)
+    %tmp4_14 = tt.expand_dims %tmp4 {axis = 1 : i32} : tensor<64xf32> -> tensor<64x1xf32> loc(#loc78)
+    %0 = tt.splat %out_ptr1 : !tt.ptr<f32> -> tensor<64x1x!tt.ptr<f32>> loc(#loc38)
+    %1 = tt.addptr %0, %xindex_9 : tensor<64x1x!tt.ptr<f32>>, tensor<64x1xi32> loc(#loc38)
+    tt.store %1, %tmp4_14 : tensor<64x1x!tt.ptr<f32>> loc(#loc39)
+    tt.return loc(#loc40)
+  } loc(#loc)
+} loc(#loc)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":33:40)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":30:19)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":29:29)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":23:28)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":23:33)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:36)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:44)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":24:23)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":26:27)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":26:37)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":28:19)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":29:21)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":34:31)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":35:29)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:45)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:41)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:55)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:50)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:68)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:60)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:34)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:73)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":39:127)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:45)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:41)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:34)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:50)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":40:104)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":41:22)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":43:23)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":44:40)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":44:8)
+#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
+#loc36 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":45:28)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:25)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:36)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/og/cogol55cthk4zevsy3dlqiyzipefv735ge2wtaddvk436qht5nox.py":49:4)
+#loc46 = loc("x2"(#loc3))
+#loc47 = loc("x1"(#loc4))
+#loc48 = loc("xoffset"(#loc5))
+#loc49 = loc("xoffset"(#loc6))
+#loc50 = loc("xindex"(#loc7))
+#loc51 = loc("xindex"(#loc8))
+#loc52 = loc("xindex"(#loc9))
+#loc53 = loc("r0_base"(#loc10))
+#loc54 = loc("r0_base"(#loc11))
+#loc55 = loc("x0"(#loc12))
+#loc56 = loc("x1"(#loc13))
+#loc57 = loc("_tmp4"(#loc2))
+#loc58 = loc("r0_index"(#loc14))
+#loc59 = loc("r0_mask"(#loc15))
+#loc60 = loc("tmp0"(#loc16))
+#loc61 = loc("tmp0"(#loc17))
+#loc62 = loc("tmp0"(#loc18))
+#loc63 = loc("tmp0"(#loc19))
+#loc64 = loc("tmp0"(#loc20))
+#loc65 = loc("tmp0"(#loc21))
+#loc66 = loc("tmp0"(#loc22))
+#loc67 = loc("tmp0"(#loc23))
+#loc68 = loc("tmp0"(#loc24))
+#loc69 = loc("tmp1"(#loc25))
+#loc70 = loc("tmp1"(#loc26))
+#loc71 = loc("tmp1"(#loc27))
+#loc72 = loc("tmp1"(#loc28))
+#loc73 = loc("tmp1"(#loc29))
+#loc74 = loc("tmp2"(#loc30))
+#loc75 = loc("tmp5"(#loc31))
+#loc76 = loc("_tmp4"(#loc32))
+#loc78 = loc("tmp4"(#loc37))
+#loc79 = loc(callsite(#loc34 at #loc77))
+#loc81 = loc(callsite(#loc36 at #loc79))

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/__grp__triton_poi_fused__to_copy_6.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_poi_fused__to_copy_6.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.source", "triton_poi_fused__to_copy_6.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.ttir", "triton_poi_fused__to_copy_6.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.ttgir", "triton_poi_fused__to_copy_6.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.llir", "triton_poi_fused__to_copy_6.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.ptx", "triton_poi_fused__to_copy_6.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.cubin", "triton_poi_fused__to_copy_6.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.json"}}

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.json

@@ -0,0 +1 @@
+{"hash": "42bb19ed265ccbd7b1943ef4d4583a35db4f47c232332144f4da3ddc335af286", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 4, "num_ctas": 1, "num_stages": 1, "warp_size": 32, "maxnreg": null, "cluster_dims": [1, 1, 1], "ptx_version": null, "ptx_options": null, "ir_override": null, "enable_fp_fusion": true, "launch_cooperative_grid": false, "launch_pdl": false, "supported_fp8_dtypes": ["fp8e4b15", "fp8e4nv", "fp8e5"], "deprecated_fp8_dot_operand_dtypes": ["fp8e4b15"], "default_dot_input_precision": "tf32", "allowed_dot_input_precisions": ["tf32", "tf32x3", "ieee"], "max_num_imprecise_acc_default": 1073741824, "extern_libs": [["libdevice", "/workspace/specforge/lib/python3.11/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]], "debug": true, "backend_name": "cuda", "sanitize_overflow": false, "arch": "sm90", "instrumentation_mode": "", "triton_version": "3.5.1", "tensordesc_meta": [], "shared": 0, "tmem_size": 0, "global_scratch_size": 0, "global_scratch_align": 1, "profile_scratch_size": 0, "profile_scratch_align": 1, "name": "triton_poi_fused__to_copy_6"}

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.llir

@@ -0,0 +1,89 @@
+; ModuleID = 'LLVMDialectModule'
+source_filename = "LLVMDialectModule"
+target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+
+; Function Attrs: nounwind
+define ptx_kernel void @triton_poi_fused__to_copy_6(ptr addrspace(1) %0, ptr addrspace(1) %1, i64 %2, i64 %3, i64 %4, i32 %5, ptr addrspace(1) readnone captures(none) %6, ptr addrspace(1) readnone captures(none) %7) local_unnamed_addr #0 !dbg !4 {
+  %9 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !7
+  %10 = shl i32 %9, 7, !dbg !8
+  %11 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !9
+  %12 = and i32 %11, 127, !dbg !9
+  %13 = or disjoint i32 %10, %12, !dbg !10
+  %14 = icmp slt i32 %13, %5, !dbg !11
+  %15 = sext i32 %13 to i64, !dbg !12
+  %.frozen = freeze i64 %2, !dbg !13
+  %16 = sdiv i64 %15, %.frozen, !dbg !13
+  %17 = mul i64 %16, %.frozen, !dbg !12
+  %.decomposed = sub i64 %15, %17, !dbg !12
+  %18 = srem i64 %16, %3, !dbg !14
+  %19 = sdiv i64 %15, %4, !dbg !15
+  %20 = insertelement <2 x i64> poison, i64 %3, i64 0, !dbg !16
+  %21 = insertelement <2 x i64> %20, i64 %2, i64 1, !dbg !16
+  %22 = icmp slt <2 x i64> %21, splat (i64 2), !dbg !16
+  %23 = icmp sgt <2 x i64> %21, splat (i64 1), !dbg !17
+  %24 = select <2 x i1> %23, <2 x i64> %21, <2 x i64> zeroinitializer, !dbg !18
+  %25 = zext <2 x i1> %22 to <2 x i64>, !dbg !19
+  %26 = add <2 x i64> %24, %25, !dbg !20
+  %27 = extractelement <2 x i64> %26, i64 0, !dbg !21
+  %28 = mul i64 %.decomposed, %27, !dbg !22
+  %29 = extractelement <2 x i64> %26, i64 1, !dbg !21
+  %30 = mul i64 %27, %29, !dbg !21
+  %31 = mul i64 %30, %19, !dbg !23
+  %32 = getelementptr i64, ptr addrspace(1) %0, i64 %18, !dbg !24
+  %33 = getelementptr i64, ptr addrspace(1) %32, i64 %28, !dbg !24
+  %34 = getelementptr i64, ptr addrspace(1) %33, i64 %31, !dbg !24
+  %35 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #2, !dbg !25
+  %36 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09@$3 ld.global.L1::evict_last.L2::cache_hint.b64 { $0 }, [ $1 + 0 ], $2;", "=l,l,l,b"(ptr addrspace(1) %34, i64 %35, i1 %14) #2, !dbg !25
+  %37 = trunc i64 %36 to i32, !dbg !26
+  %38 = mul i64 %18, %29, !dbg !27
+  %39 = getelementptr i32, ptr addrspace(1) %1, i64 %.decomposed, !dbg !28
+  %40 = getelementptr i32, ptr addrspace(1) %39, i64 %38, !dbg !28
+  %41 = getelementptr i32, ptr addrspace(1) %40, i64 %31, !dbg !28
+  tail call void asm sideeffect "@$2 st.global.b32 [ $1 + 0 ], { $0 };", "r,l,b"(i32 %37, ptr addrspace(1) %41, i1 %14) #2, !dbg !29
+  ret void, !dbg !30
+}
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
+declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
+
+attributes #0 = { nounwind "nvvm.reqntid"="128" }
+attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+attributes #2 = { nounwind }
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
+!1 = !DIFile(filename: "c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35")
+!2 = !{i32 2, !"Debug Info Version", i32 3}
+!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
+!4 = distinct !DISubprogram(name: "triton_poi_fused__to_copy_6", linkageName: "triton_poi_fused__to_copy_6", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
+!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
+!6 = !{}
+!7 = !DILocation(line: 19, column: 28, scope: !4)
+!8 = !DILocation(line: 19, column: 33, scope: !4)
+!9 = !DILocation(line: 20, column: 36, scope: !4)
+!10 = !DILocation(line: 20, column: 23, scope: !4)
+!11 = !DILocation(line: 21, column: 21, scope: !4)
+!12 = !DILocation(line: 22, column: 19, scope: !4)
+!13 = !DILocation(line: 23, column: 21, scope: !4)
+!14 = !DILocation(line: 23, column: 28, scope: !4)
+!15 = !DILocation(line: 24, column: 19, scope: !4)
+!16 = !DILocation(line: 25, column: 54, scope: !4)
+!17 = !DILocation(line: 25, column: 80, scope: !4)
+!18 = !DILocation(line: 25, column: 71, scope: !4)
+!19 = !DILocation(line: 25, scope: !4)
+!20 = !DILocation(line: 25, column: 62, scope: !4)
+!21 = !DILocation(line: 25, column: 91, scope: !4)
+!22 = !DILocation(line: 25, column: 39, scope: !4)
+!23 = !DILocation(line: 25, column: 138, scope: !4)
+!24 = !DILocation(line: 25, column: 30, scope: !4)
+!25 = !DILocation(line: 25, column: 186, scope: !4)
+!26 = !DILocation(line: 26, column: 19, scope: !4)
+!27 = !DILocation(line: 27, column: 34, scope: !4)
+!28 = !DILocation(line: 27, column: 25, scope: !4)
+!29 = !DILocation(line: 27, column: 187, scope: !4)
+!30 = !DILocation(line: 27, column: 4, scope: !4)

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.ptx

@@ -0,0 +1,311 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.7
+.target sm_90a
+.address_size 64
+
+	// .globl	triton_poi_fused__to_copy_6 // -- Begin function triton_poi_fused__to_copy_6
+                                        // @triton_poi_fused__to_copy_6
+.visible .entry triton_poi_fused__to_copy_6(
+	.param .u64 .ptr .global .align 1 triton_poi_fused__to_copy_6_param_0,
+	.param .u64 .ptr .global .align 1 triton_poi_fused__to_copy_6_param_1,
+	.param .u64 triton_poi_fused__to_copy_6_param_2,
+	.param .u64 triton_poi_fused__to_copy_6_param_3,
+	.param .u64 triton_poi_fused__to_copy_6_param_4,
+	.param .u32 triton_poi_fused__to_copy_6_param_5,
+	.param .u64 .ptr .global .align 1 triton_poi_fused__to_copy_6_param_6,
+	.param .u64 .ptr .global .align 1 triton_poi_fused__to_copy_6_param_7
+)
+.reqntid 128
+{
+	.reg .pred 	%p<10>;
+	.reg .b32 	%r<20>;
+	.reg .b64 	%rd<54>;
+	.loc	1 18 0                          // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:18:0
+$L__func_begin0:
+	.loc	1 18 0                          // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:18:0
+
+// %bb.0:
+	ld.param.b64 	%rd16, [triton_poi_fused__to_copy_6_param_3];
+	ld.param.b64 	%rd15, [triton_poi_fused__to_copy_6_param_2];
+$L__tmp0:
+	.loc	1 19 28                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:19:28
+	mov.u32 	%r2, %ctaid.x;
+	.loc	1 19 33                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:19:33
+	shl.b32 	%r3, %r2, 7;
+	.loc	1 20 36                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:20:36
+	mov.u32 	%r4, %tid.x;
+	and.b32 	%r5, %r4, 127;
+	.loc	1 20 23                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:20:23
+	or.b32 	%r6, %r3, %r5;
+	.loc	1 22 19                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:22:19
+	cvt.s64.s32 	%rd1, %r6;
+	.loc	1 23 21                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:23:21
+	or.b64 	%rd19, %rd1, %rd15;
+	and.b64 	%rd20, %rd19, -4294967296;
+	setp.ne.b64 	%p1, %rd20, 0;
+	cvt.u32.u64 	%r19, %rd1;
+	@%p1 bra 	$L__BB0_2;
+	bra.uni 	$L__BB0_1;
+$L__BB0_2:
+	div.s64 	%rd51, %rd1, %rd15;
+	bra.uni 	$L__BB0_3;
+$L__BB0_1:
+	cvt.u32.u64 	%r7, %rd15;
+	div.u32 	%r9, %r19, %r7;
+	cvt.u64.u32 	%rd51, %r9;
+$L__BB0_3:
+	.loc	1 0 21                          // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:0:21
+	ld.param.b64 	%rd17, [triton_poi_fused__to_copy_6_param_4];
+	.loc	1 22 19                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:22:19
+	mul.lo.s64 	%rd21, %rd51, %rd15;
+	.loc	1 23 28                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:23:28
+	or.b64 	%rd22, %rd51, %rd16;
+	and.b64 	%rd23, %rd22, -4294967296;
+	setp.ne.b64 	%p2, %rd23, 0;
+	@%p2 bra 	$L__BB0_5;
+	bra.uni 	$L__BB0_4;
+$L__BB0_5:
+	rem.s64 	%rd52, %rd51, %rd16;
+	bra.uni 	$L__BB0_6;
+$L__BB0_4:
+	cvt.u32.u64 	%r10, %rd16;
+	cvt.u32.u64 	%r11, %rd51;
+	rem.u32 	%r12, %r11, %r10;
+	cvt.u64.u32 	%rd52, %r12;
+$L__BB0_6:
+	.loc	1 0 28                          // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:0:28
+	ld.param.b32 	%r1, [triton_poi_fused__to_copy_6_param_5];
+	ld.param.b64 	%rd14, [triton_poi_fused__to_copy_6_param_1];
+	ld.param.b64 	%rd13, [triton_poi_fused__to_copy_6_param_0];
+	sub.s64 	%rd6, %rd1, %rd21;
+	.loc	1 24 19                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:24:19
+	or.b64 	%rd24, %rd1, %rd17;
+	and.b64 	%rd25, %rd24, -4294967296;
+	setp.ne.b64 	%p3, %rd25, 0;
+	@%p3 bra 	$L__BB0_8;
+	bra.uni 	$L__BB0_7;
+$L__BB0_8:
+	div.s64 	%rd53, %rd1, %rd17;
+	bra.uni 	$L__BB0_9;
+$L__BB0_7:
+	cvt.u32.u64 	%r13, %rd17;
+	div.u32 	%r15, %r19, %r13;
+	cvt.u64.u32 	%rd53, %r15;
+$L__BB0_9:
+	.loc	1 21 21                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:21:21
+	setp.lt.s32 	%p4, %r19, %r1;
+	.loc	1 25 54                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:54
+	setp.lt.s64 	%p6, %rd15, 2;
+	setp.lt.s64 	%p7, %rd16, 2;
+	.loc	1 25 80                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:80
+	setp.gt.s64 	%p8, %rd15, 1;
+	setp.gt.s64 	%p9, %rd16, 1;
+	.loc	1 25 71                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:71
+	selp.b64 	%rd31, %rd16, 0, %p9;
+	selp.b64 	%rd32, %rd15, 0, %p8;
+	.loc	1 25 0                          // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25
+	selp.b64 	%rd33, 1, 0, %p7;
+	selp.b64 	%rd34, 1, 0, %p6;
+	.loc	1 25 62                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:62
+	add.s64 	%rd35, %rd32, %rd34;
+	add.s64 	%rd36, %rd31, %rd33;
+	.loc	1 25 39                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:39
+	mul.lo.s64 	%rd37, %rd6, %rd36;
+	.loc	1 25 91                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:91
+	mul.lo.s64 	%rd38, %rd36, %rd35;
+	.loc	1 25 138                        // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:138
+	mul.lo.s64 	%rd39, %rd38, %rd53;
+	.loc	1 25 30                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:30
+	shl.b64 	%rd40, %rd52, 3;
+	add.s64 	%rd41, %rd13, %rd40;
+	shl.b64 	%rd42, %rd37, 3;
+	add.s64 	%rd43, %rd41, %rd42;
+	shl.b64 	%rd44, %rd39, 3;
+	add.s64 	%rd28, %rd43, %rd44;
+	.loc	1 25 186                        // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:25:186
+	// begin inline asm
+	mov.u64 %rd29, 0x0;
+	createpolicy.fractional.L2::evict_last.b64 %rd29, 1.0;
+	// end inline asm
+	// begin inline asm
+	mov.u64 %rd27, 0x0;
+	@%p4 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd27 }, [ %rd28 + 0 ], %rd29;
+	// end inline asm
+	.loc	1 26 19                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:26:19
+	cvt.u32.u64 	%r16, %rd27;
+	.loc	1 27 34                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:27:34
+	mul.lo.s64 	%rd45, %rd52, %rd35;
+	.loc	1 27 25                         // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:27:25
+	shl.b64 	%rd46, %rd6, 2;
+	add.s64 	%rd47, %rd14, %rd46;
+	shl.b64 	%rd48, %rd45, 2;
+	add.s64 	%rd49, %rd47, %rd48;
+	shl.b64 	%rd50, %rd39, 2;
+	add.s64 	%rd30, %rd49, %rd50;
+	.loc	1 27 187                        // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:27:187
+	// begin inline asm
+	@%p4 st.global.b32 [ %rd30 + 0 ], { %r16 };
+	// end inline asm
+	.loc	1 27 4                          // c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py:27:4
+	ret;
+$L__tmp1:
+$L__func_end0:
+                                        // -- End function
+}
+	.file	1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py"
+	.section	.debug_abbrev
+	{
+.b8 1                                   // Abbreviation Code
+.b8 17                                  // DW_TAG_compile_unit
+.b8 0                                   // DW_CHILDREN_no
+.b8 37                                  // DW_AT_producer
+.b8 8                                   // DW_FORM_string
+.b8 19                                  // DW_AT_language
+.b8 5                                   // DW_FORM_data2
+.b8 3                                   // DW_AT_name
+.b8 8                                   // DW_FORM_string
+.b8 16                                  // DW_AT_stmt_list
+.b8 6                                   // DW_FORM_data4
+.b8 27                                  // DW_AT_comp_dir
+.b8 8                                   // DW_FORM_string
+.b8 0                                   // EOM(1)
+.b8 0                                   // EOM(2)
+.b8 0                                   // EOM(3)
+	}
+	.section	.debug_info
+	{
+.b32 135                                // Length of Unit
+.b8 2                                   // DWARF version number
+.b8 0
+.b32 .debug_abbrev                      // Offset Into Abbrev. Section
+.b8 8                                   // Address Size (in bytes)
+.b8 1                                   // Abbrev [1] 0xb:0x80 DW_TAG_compile_unit
+.b8 116                                 // DW_AT_producer
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2                                   // DW_AT_language
+.b8 0
+.b8 99                                  // DW_AT_name
+.b8 51
+.b8 53
+.b8 104
+.b8 117
+.b8 113
+.b8 112
+.b8 54
+.b8 110
+.b8 103
+.b8 122
+.b8 104
+.b8 54
+.b8 55
+.b8 107
+.b8 116
+.b8 51
+.b8 50
+.b8 107
+.b8 117
+.b8 120
+.b8 111
+.b8 113
+.b8 112
+.b8 103
+.b8 104
+.b8 99
+.b8 51
+.b8 50
+.b8 102
+.b8 115
+.b8 116
+.b8 118
+.b8 52
+.b8 122
+.b8 111
+.b8 103
+.b8 99
+.b8 111
+.b8 117
+.b8 122
+.b8 97
+.b8 98
+.b8 100
+.b8 120
+.b8 120
+.b8 119
+.b8 116
+.b8 97
+.b8 51
+.b8 115
+.b8 108
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line                        // DW_AT_stmt_list
+.b8 47                                  // DW_AT_comp_dir
+.b8 119
+.b8 111
+.b8 114
+.b8 107
+.b8 115
+.b8 112
+.b8 97
+.b8 99
+.b8 101
+.b8 47
+.b8 104
+.b8 97
+.b8 110
+.b8 114
+.b8 117
+.b8 105
+.b8 47
+.b8 83
+.b8 112
+.b8 101
+.b8 99
+.b8 70
+.b8 111
+.b8 114
+.b8 103
+.b8 101
+.b8 45
+.b8 101
+.b8 120
+.b8 116
+.b8 47
+.b8 99
+.b8 97
+.b8 99
+.b8 104
+.b8 101
+.b8 47
+.b8 99
+.b8 111
+.b8 109
+.b8 112
+.b8 105
+.b8 108
+.b8 101
+.b8 100
+.b8 95
+.b8 107
+.b8 101
+.b8 114
+.b8 110
+.b8 101
+.b8 108
+.b8 115
+.b8 47
+.b8 51
+.b8 53
+.b8 0
+	}
+	.section	.debug_macinfo	{	}

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.source

@@ -0,0 +1,226 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":18:0)
+#loc56 = loc("in_ptr0"(#loc))
+#loc57 = loc("out_ptr0"(#loc))
+#loc58 = loc("ks0"(#loc))
+#loc59 = loc("ks1"(#loc))
+#loc60 = loc("ks2"(#loc))
+#loc61 = loc("xnumel"(#loc))
+module {
+  tt.func public @triton_poi_fused__to_copy_6(%in_ptr0: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %ks2: i64 loc("ks2"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
+    %xoffset = tt.get_program_id x : i32 loc(#loc62)
+    %xoffset_0 = arith.constant 128 : i32 loc(#loc63)
+    %xoffset_1 = arith.constant 128 : i32 loc(#loc63)
+    %xoffset_2 = arith.muli %xoffset, %xoffset_1 : i32 loc(#loc63)
+    %xindex = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc64)
+    %xindex_3 = tt.splat %xoffset_2 : i32 -> tensor<128xi32> loc(#loc65)
+    %xindex_4 = arith.addi %xindex_3, %xindex : tensor<128xi32> loc(#loc65)
+    %xmask = tt.splat %xnumel : i32 -> tensor<128xi32> loc(#loc66)
+    %xmask_5 = arith.cmpi slt, %xindex_4, %xmask : tensor<128xi32> loc(#loc66)
+    %x0 = arith.extsi %xindex_4 : tensor<128xi32> to tensor<128xi64> loc(#loc67)
+    %x0_6 = tt.splat %ks0 : i64 -> tensor<128xi64> loc(#loc67)
+    %x0_7 = arith.remsi %x0, %x0_6 : tensor<128xi64> loc(#loc67)
+    %x1 = arith.extsi %xindex_4 : tensor<128xi32> to tensor<128xi64> loc(#loc68)
+    %x1_8 = tt.splat %ks0 : i64 -> tensor<128xi64> loc(#loc68)
+    %x1_9 = arith.divsi %x1, %x1_8 : tensor<128xi64> loc(#loc68)
+    %x1_10 = tt.splat %ks1 : i64 -> tensor<128xi64> loc(#loc69)
+    %x1_11 = arith.remsi %x1_9, %x1_10 : tensor<128xi64> loc(#loc69)
+    %x2 = arith.extsi %xindex_4 : tensor<128xi32> to tensor<128xi64> loc(#loc70)
+    %x2_12 = tt.splat %ks2 : i64 -> tensor<128xi64> loc(#loc70)
+    %x2_13 = arith.divsi %x2, %x2_12 : tensor<128xi64> loc(#loc70)
+    %tmp0 = arith.constant 1 : i32 loc(#loc71)
+    %tmp0_14 = arith.extsi %tmp0 : i32 to i64 loc(#loc71)
+    %tmp0_15 = arith.cmpi sge, %tmp0_14, %ks1 : i64 loc(#loc71)
+    %tmp0_16 = arith.constant 1 : i32 loc(#loc72)
+    %tmp0_17 = arith.constant 1 : i32 loc(#loc72)
+    %tmp0_18 = arith.extui %tmp0_15 : i1 to i32 loc(#loc72)
+    %tmp0_19 = arith.muli %tmp0_17, %tmp0_18 : i32 loc(#loc72)
+    %tmp0_20 = arith.constant 1 : i32 loc(#loc73)
+    %tmp0_21 = arith.extsi %tmp0_20 : i32 to i64 loc(#loc73)
+    %tmp0_22 = arith.cmpi sgt, %ks1, %tmp0_21 : i64 loc(#loc73)
+    %tmp0_23 = arith.extui %tmp0_22 : i1 to i64 loc(#loc74)
+    %tmp0_24 = arith.muli %ks1, %tmp0_23 : i64 loc(#loc74)
+    %tmp0_25 = arith.extsi %tmp0_19 : i32 to i64 loc(#loc75)
+    %tmp0_26 = arith.addi %tmp0_25, %tmp0_24 : i64 loc(#loc75)
+    %tmp0_27 = tt.splat %tmp0_26 : i64 -> tensor<128xi64> loc(#loc76)
+    %tmp0_28 = arith.muli %x0_7, %tmp0_27 : tensor<128xi64> loc(#loc76)
+    %tmp0_29 = arith.addi %x1_11, %tmp0_28 : tensor<128xi64> loc(#loc77)
+    %tmp0_30 = arith.constant 1 : i32 loc(#loc78)
+    %tmp0_31 = arith.extsi %tmp0_30 : i32 to i64 loc(#loc78)
+    %tmp0_32 = arith.cmpi sge, %tmp0_31, %ks0 : i64 loc(#loc78)
+    %tmp0_33 = arith.constant 1 : i32 loc(#loc79)
+    %tmp0_34 = arith.constant 1 : i32 loc(#loc79)
+    %tmp0_35 = arith.extui %tmp0_32 : i1 to i32 loc(#loc79)
+    %tmp0_36 = arith.muli %tmp0_34, %tmp0_35 : i32 loc(#loc79)
+    %tmp0_37 = arith.constant 1 : i32 loc(#loc80)
+    %tmp0_38 = arith.extsi %tmp0_37 : i32 to i64 loc(#loc80)
+    %tmp0_39 = arith.cmpi sgt, %ks0, %tmp0_38 : i64 loc(#loc80)
+    %tmp0_40 = arith.extui %tmp0_39 : i1 to i64 loc(#loc81)
+    %tmp0_41 = arith.muli %ks0, %tmp0_40 : i64 loc(#loc81)
+    %tmp0_42 = arith.extsi %tmp0_36 : i32 to i64 loc(#loc82)
+    %tmp0_43 = arith.addi %tmp0_42, %tmp0_41 : i64 loc(#loc82)
+    %tmp0_44 = tt.splat %tmp0_43 : i64 -> tensor<128xi64> loc(#loc83)
+    %tmp0_45 = arith.muli %x2_13, %tmp0_44 : tensor<128xi64> loc(#loc83)
+    %tmp0_46 = arith.constant 1 : i32 loc(#loc84)
+    %tmp0_47 = arith.extsi %tmp0_46 : i32 to i64 loc(#loc84)
+    %tmp0_48 = arith.cmpi sge, %tmp0_47, %ks1 : i64 loc(#loc84)
+    %tmp0_49 = arith.constant 1 : i32 loc(#loc85)
+    %tmp0_50 = arith.constant 1 : i32 loc(#loc85)
+    %tmp0_51 = arith.extui %tmp0_48 : i1 to i32 loc(#loc85)
+    %tmp0_52 = arith.muli %tmp0_50, %tmp0_51 : i32 loc(#loc85)
+    %tmp0_53 = arith.constant 1 : i32 loc(#loc86)
+    %tmp0_54 = arith.extsi %tmp0_53 : i32 to i64 loc(#loc86)
+    %tmp0_55 = arith.cmpi sgt, %ks1, %tmp0_54 : i64 loc(#loc86)
+    %tmp0_56 = arith.extui %tmp0_55 : i1 to i64 loc(#loc87)
+    %tmp0_57 = arith.muli %ks1, %tmp0_56 : i64 loc(#loc87)
+    %tmp0_58 = arith.extsi %tmp0_52 : i32 to i64 loc(#loc88)
+    %tmp0_59 = arith.addi %tmp0_58, %tmp0_57 : i64 loc(#loc88)
+    %tmp0_60 = tt.splat %tmp0_59 : i64 -> tensor<128xi64> loc(#loc89)
+    %tmp0_61 = arith.muli %tmp0_45, %tmp0_60 : tensor<128xi64> loc(#loc89)
+    %tmp0_62 = arith.addi %tmp0_29, %tmp0_61 : tensor<128xi64> loc(#loc90)
+    %tmp0_63 = tt.splat %in_ptr0 : !tt.ptr<i64> -> tensor<128x!tt.ptr<i64>> loc(#loc91)
+    %tmp0_64 = tt.addptr %tmp0_63, %tmp0_62 : tensor<128x!tt.ptr<i64>>, tensor<128xi64> loc(#loc91)
+    %tmp0_65 = tt.load %tmp0_64, %xmask_5 evictionPolicy = evict_last : tensor<128x!tt.ptr<i64>> loc(#loc92)
+    %tmp1 = arith.trunci %tmp0_65 : tensor<128xi64> to tensor<128xi32> loc(#loc93)
+    %c1_i32 = arith.constant 1 : i32 loc(#loc33)
+    %0 = arith.extsi %c1_i32 : i32 to i64 loc(#loc33)
+    %1 = arith.cmpi sge, %0, %ks0 : i64 loc(#loc33)
+    %c1_i32_66 = arith.constant 1 : i32 loc(#loc34)
+    %c1_i32_67 = arith.constant 1 : i32 loc(#loc34)
+    %2 = arith.extui %1 : i1 to i32 loc(#loc34)
+    %3 = arith.muli %c1_i32_67, %2 : i32 loc(#loc34)
+    %c1_i32_68 = arith.constant 1 : i32 loc(#loc35)
+    %4 = arith.extsi %c1_i32_68 : i32 to i64 loc(#loc35)
+    %5 = arith.cmpi sgt, %ks0, %4 : i64 loc(#loc35)
+    %6 = arith.extui %5 : i1 to i64 loc(#loc36)
+    %7 = arith.muli %ks0, %6 : i64 loc(#loc36)
+    %8 = arith.extsi %3 : i32 to i64 loc(#loc37)
+    %9 = arith.addi %8, %7 : i64 loc(#loc37)
+    %10 = tt.splat %9 : i64 -> tensor<128xi64> loc(#loc38)
+    %11 = arith.muli %x1_11, %10 : tensor<128xi64> loc(#loc38)
+    %12 = arith.addi %x0_7, %11 : tensor<128xi64> loc(#loc39)
+    %c1_i32_69 = arith.constant 1 : i32 loc(#loc40)
+    %13 = arith.extsi %c1_i32_69 : i32 to i64 loc(#loc40)
+    %14 = arith.cmpi sge, %13, %ks0 : i64 loc(#loc40)
+    %c1_i32_70 = arith.constant 1 : i32 loc(#loc41)
+    %c1_i32_71 = arith.constant 1 : i32 loc(#loc41)
+    %15 = arith.extui %14 : i1 to i32 loc(#loc41)
+    %16 = arith.muli %c1_i32_71, %15 : i32 loc(#loc41)
+    %c1_i32_72 = arith.constant 1 : i32 loc(#loc42)
+    %17 = arith.extsi %c1_i32_72 : i32 to i64 loc(#loc42)
+    %18 = arith.cmpi sgt, %ks0, %17 : i64 loc(#loc42)
+    %19 = arith.extui %18 : i1 to i64 loc(#loc43)
+    %20 = arith.muli %ks0, %19 : i64 loc(#loc43)
+    %21 = arith.extsi %16 : i32 to i64 loc(#loc44)
+    %22 = arith.addi %21, %20 : i64 loc(#loc44)
+    %23 = tt.splat %22 : i64 -> tensor<128xi64> loc(#loc45)
+    %24 = arith.muli %x2_13, %23 : tensor<128xi64> loc(#loc45)
+    %c1_i32_73 = arith.constant 1 : i32 loc(#loc46)
+    %25 = arith.extsi %c1_i32_73 : i32 to i64 loc(#loc46)
+    %26 = arith.cmpi sge, %25, %ks1 : i64 loc(#loc46)
+    %c1_i32_74 = arith.constant 1 : i32 loc(#loc47)
+    %c1_i32_75 = arith.constant 1 : i32 loc(#loc47)
+    %27 = arith.extui %26 : i1 to i32 loc(#loc47)
+    %28 = arith.muli %c1_i32_75, %27 : i32 loc(#loc47)
+    %c1_i32_76 = arith.constant 1 : i32 loc(#loc48)
+    %29 = arith.extsi %c1_i32_76 : i32 to i64 loc(#loc48)
+    %30 = arith.cmpi sgt, %ks1, %29 : i64 loc(#loc48)
+    %31 = arith.extui %30 : i1 to i64 loc(#loc49)
+    %32 = arith.muli %ks1, %31 : i64 loc(#loc49)
+    %33 = arith.extsi %28 : i32 to i64 loc(#loc50)
+    %34 = arith.addi %33, %32 : i64 loc(#loc50)
+    %35 = tt.splat %34 : i64 -> tensor<128xi64> loc(#loc51)
+    %36 = arith.muli %24, %35 : tensor<128xi64> loc(#loc51)
+    %37 = arith.addi %12, %36 : tensor<128xi64> loc(#loc52)
+    %38 = tt.splat %out_ptr0 : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>> loc(#loc53)
+    %39 = tt.addptr %38, %37 : tensor<128x!tt.ptr<i32>>, tensor<128xi64> loc(#loc53)
+    tt.store %39, %tmp1, %xmask_5 : tensor<128x!tt.ptr<i32>> loc(#loc54)
+    tt.return loc(#loc55)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":19:28)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":19:33)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":20:36)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":20:23)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":21:21)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":22:19)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":23:21)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":23:28)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":24:19)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:54)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:46)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:80)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:71)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:62)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:39)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:35)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:106)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:98)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:132)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:123)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:114)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:91)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:153)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:145)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:179)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:170)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:161)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:138)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:87)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:30)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:186)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":26:19)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:49)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:41)
+#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:75)
+#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:66)
+#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:57)
+#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:34)
+#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:30)
+#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:101)
+#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:93)
+#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:127)
+#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:118)
+#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:109)
+#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:86)
+#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:148)
+#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:140)
+#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:174)
+#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:165)
+#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:156)
+#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:133)
+#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:82)
+#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:25)
+#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:187)
+#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:4)
+#loc62 = loc("xoffset"(#loc1))
+#loc63 = loc("xoffset"(#loc2))
+#loc64 = loc("xindex"(#loc3))
+#loc65 = loc("xindex"(#loc4))
+#loc66 = loc("xmask"(#loc5))
+#loc67 = loc("x0"(#loc6))
+#loc68 = loc("x1"(#loc7))
+#loc69 = loc("x1"(#loc8))
+#loc70 = loc("x2"(#loc9))
+#loc71 = loc("tmp0"(#loc10))
+#loc72 = loc("tmp0"(#loc11))
+#loc73 = loc("tmp0"(#loc12))
+#loc74 = loc("tmp0"(#loc13))
+#loc75 = loc("tmp0"(#loc14))
+#loc76 = loc("tmp0"(#loc15))
+#loc77 = loc("tmp0"(#loc16))
+#loc78 = loc("tmp0"(#loc17))
+#loc79 = loc("tmp0"(#loc18))
+#loc80 = loc("tmp0"(#loc19))
+#loc81 = loc("tmp0"(#loc20))
+#loc82 = loc("tmp0"(#loc21))
+#loc83 = loc("tmp0"(#loc22))
+#loc84 = loc("tmp0"(#loc23))
+#loc85 = loc("tmp0"(#loc24))
+#loc86 = loc("tmp0"(#loc25))
+#loc87 = loc("tmp0"(#loc26))
+#loc88 = loc("tmp0"(#loc27))
+#loc89 = loc("tmp0"(#loc28))
+#loc90 = loc("tmp0"(#loc29))
+#loc91 = loc("tmp0"(#loc30))
+#loc92 = loc("tmp0"(#loc31))
+#loc93 = loc("tmp1"(#loc32))

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.ttgir

@@ -0,0 +1,122 @@
+#blocked = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":18:0)
+#loc35 = loc("in_ptr0"(#loc))
+#loc36 = loc("out_ptr0"(#loc))
+#loc37 = loc("ks0"(#loc))
+#loc38 = loc("ks1"(#loc))
+#loc39 = loc("ks2"(#loc))
+#loc40 = loc("xnumel"(#loc))
+module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
+  tt.func public @triton_poi_fused__to_copy_6(%in_ptr0: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %ks2: i64 loc("ks2"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
+    %c128_i32 = arith.constant 128 : i32 loc(#loc1)
+    %c1_i64 = arith.constant 1 : i64 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc41)
+    %xoffset_0 = arith.muli %xoffset, %c128_i32 : i32 loc(#loc42)
+    %xindex = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #blocked> loc(#loc43)
+    %xindex_1 = tt.splat %xoffset_0 : i32 -> tensor<128xi32, #blocked> loc(#loc44)
+    %xindex_2 = arith.addi %xindex_1, %xindex : tensor<128xi32, #blocked> loc(#loc44)
+    %xmask = tt.splat %xnumel : i32 -> tensor<128xi32, #blocked> loc(#loc45)
+    %xmask_3 = arith.cmpi slt, %xindex_2, %xmask : tensor<128xi32, #blocked> loc(#loc45)
+    %x0 = arith.extsi %xindex_2 : tensor<128xi32, #blocked> to tensor<128xi64, #blocked> loc(#loc46)
+    %x0_4 = tt.splat %ks0 : i64 -> tensor<128xi64, #blocked> loc(#loc46)
+    %x0_5 = arith.remsi %x0, %x0_4 : tensor<128xi64, #blocked> loc(#loc46)
+    %x1 = arith.divsi %x0, %x0_4 : tensor<128xi64, #blocked> loc(#loc47)
+    %x1_6 = tt.splat %ks1 : i64 -> tensor<128xi64, #blocked> loc(#loc48)
+    %x1_7 = arith.remsi %x1, %x1_6 : tensor<128xi64, #blocked> loc(#loc48)
+    %x2 = tt.splat %ks2 : i64 -> tensor<128xi64, #blocked> loc(#loc49)
+    %x2_8 = arith.divsi %x0, %x2 : tensor<128xi64, #blocked> loc(#loc49)
+    %tmp0 = arith.cmpi sle, %ks1, %c1_i64 : i64 loc(#loc50)
+    %tmp0_9 = arith.cmpi sgt, %ks1, %c1_i64 : i64 loc(#loc51)
+    %tmp0_10 = arith.extui %tmp0_9 : i1 to i64 loc(#loc52)
+    %tmp0_11 = arith.muli %ks1, %tmp0_10 : i64 loc(#loc52)
+    %tmp0_12 = arith.extui %tmp0 : i1 to i64 loc(#loc68)
+    %tmp0_13 = arith.addi %tmp0_12, %tmp0_11 : i64 loc(#loc53)
+    %tmp0_14 = tt.splat %tmp0_13 : i64 -> tensor<128xi64, #blocked> loc(#loc55)
+    %tmp0_15 = arith.muli %x0_5, %tmp0_14 : tensor<128xi64, #blocked> loc(#loc55)
+    %tmp0_16 = arith.addi %x1_7, %tmp0_15 : tensor<128xi64, #blocked> loc(#loc56)
+    %tmp0_17 = arith.cmpi sle, %ks0, %c1_i64 : i64 loc(#loc57)
+    %tmp0_18 = arith.cmpi sgt, %ks0, %c1_i64 : i64 loc(#loc58)
+    %tmp0_19 = arith.extui %tmp0_18 : i1 to i64 loc(#loc59)
+    %tmp0_20 = arith.muli %ks0, %tmp0_19 : i64 loc(#loc59)
+    %tmp0_21 = arith.extui %tmp0_17 : i1 to i64 loc(#loc69)
+    %tmp0_22 = arith.addi %tmp0_21, %tmp0_20 : i64 loc(#loc60)
+    %tmp0_23 = tt.splat %tmp0_22 : i64 -> tensor<128xi64, #blocked> loc(#loc62)
+    %tmp0_24 = arith.muli %x2_8, %tmp0_23 : tensor<128xi64, #blocked> loc(#loc62)
+    %tmp0_25 = arith.muli %tmp0_24, %tmp0_14 : tensor<128xi64, #blocked> loc(#loc63)
+    %tmp0_26 = arith.addi %tmp0_16, %tmp0_25 : tensor<128xi64, #blocked> loc(#loc64)
+    %tmp0_27 = tt.splat %in_ptr0 : !tt.ptr<i64> -> tensor<128x!tt.ptr<i64>, #blocked> loc(#loc65)
+    %tmp0_28 = tt.addptr %tmp0_27, %tmp0_26 : tensor<128x!tt.ptr<i64>, #blocked>, tensor<128xi64, #blocked> loc(#loc65)
+    %tmp0_29 = tt.load %tmp0_28, %xmask_3 evictionPolicy = evict_last : tensor<128x!tt.ptr<i64>, #blocked> loc(#loc66)
+    %tmp1 = arith.trunci %tmp0_29 : tensor<128xi64, #blocked> to tensor<128xi32, #blocked> loc(#loc67)
+    %0 = arith.muli %x1_7, %tmp0_23 : tensor<128xi64, #blocked> loc(#loc29)
+    %1 = arith.addi %x0_5, %0 : tensor<128xi64, #blocked> loc(#loc30)
+    %2 = arith.addi %1, %tmp0_25 : tensor<128xi64, #blocked> loc(#loc31)
+    %3 = tt.splat %out_ptr0 : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>, #blocked> loc(#loc32)
+    %4 = tt.addptr %3, %2 : tensor<128x!tt.ptr<i32>, #blocked>, tensor<128xi64, #blocked> loc(#loc32)
+    tt.store %4, %tmp1, %xmask_3 : tensor<128x!tt.ptr<i32>, #blocked> loc(#loc33)
+    tt.return loc(#loc34)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc(unknown)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":19:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":19:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":20:36)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":20:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":21:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":22:19)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":23:21)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":23:28)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":24:19)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:54)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:80)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:71)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:62)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:46)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:39)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:35)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:106)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:132)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:123)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:114)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:98)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:91)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:138)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:87)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:30)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:186)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":26:19)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:34)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:30)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:82)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:25)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:187)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:4)
+#loc41 = loc("xoffset"(#loc2))
+#loc42 = loc("xoffset"(#loc3))
+#loc43 = loc("xindex"(#loc4))
+#loc44 = loc("xindex"(#loc5))
+#loc45 = loc("xmask"(#loc6))
+#loc46 = loc("x0"(#loc7))
+#loc47 = loc("x1"(#loc8))
+#loc48 = loc("x1"(#loc9))
+#loc49 = loc("x2"(#loc10))
+#loc50 = loc("tmp0"(#loc11))
+#loc51 = loc("tmp0"(#loc12))
+#loc52 = loc("tmp0"(#loc13))
+#loc53 = loc("tmp0"(#loc14))
+#loc54 = loc("tmp0"(#loc15))
+#loc55 = loc("tmp0"(#loc16))
+#loc56 = loc("tmp0"(#loc17))
+#loc57 = loc("tmp0"(#loc18))
+#loc58 = loc("tmp0"(#loc19))
+#loc59 = loc("tmp0"(#loc20))
+#loc60 = loc("tmp0"(#loc21))
+#loc61 = loc("tmp0"(#loc22))
+#loc62 = loc("tmp0"(#loc23))
+#loc63 = loc("tmp0"(#loc24))
+#loc64 = loc("tmp0"(#loc25))
+#loc65 = loc("tmp0"(#loc26))
+#loc66 = loc("tmp0"(#loc27))
+#loc67 = loc("tmp1"(#loc28))
+#loc68 = loc(fused[#loc53, #loc54])
+#loc69 = loc(fused[#loc60, #loc61])

SpecForge-ext/cache/compiled_kernels/triton/6/IK5RT3JGLTF5PMMUH32NIWB2GXNU6R6CGIZSCRHU3I65YM226KDA/triton_poi_fused__to_copy_6.ttir

@@ -0,0 +1,121 @@
+#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":18:0)
+#loc35 = loc("in_ptr0"(#loc))
+#loc36 = loc("out_ptr0"(#loc))
+#loc37 = loc("ks0"(#loc))
+#loc38 = loc("ks1"(#loc))
+#loc39 = loc("ks2"(#loc))
+#loc40 = loc("xnumel"(#loc))
+module {
+  tt.func public @triton_poi_fused__to_copy_6(%in_ptr0: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %out_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %ks0: i64 loc("ks0"(#loc)), %ks1: i64 loc("ks1"(#loc)), %ks2: i64 loc("ks2"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
+    %c1_i64 = arith.constant 1 : i64 loc(#loc1)
+    %c128_i32 = arith.constant 128 : i32 loc(#loc1)
+    %xoffset = tt.get_program_id x : i32 loc(#loc41)
+    %xoffset_0 = arith.muli %xoffset, %c128_i32 : i32 loc(#loc42)
+    %xindex = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc43)
+    %xindex_1 = tt.splat %xoffset_0 : i32 -> tensor<128xi32> loc(#loc44)
+    %xindex_2 = arith.addi %xindex_1, %xindex : tensor<128xi32> loc(#loc44)
+    %xmask = tt.splat %xnumel : i32 -> tensor<128xi32> loc(#loc45)
+    %xmask_3 = arith.cmpi slt, %xindex_2, %xmask : tensor<128xi32> loc(#loc45)
+    %x0 = arith.extsi %xindex_2 : tensor<128xi32> to tensor<128xi64> loc(#loc46)
+    %x0_4 = tt.splat %ks0 : i64 -> tensor<128xi64> loc(#loc46)
+    %x0_5 = arith.remsi %x0, %x0_4 : tensor<128xi64> loc(#loc46)
+    %x1 = arith.divsi %x0, %x0_4 : tensor<128xi64> loc(#loc47)
+    %x1_6 = tt.splat %ks1 : i64 -> tensor<128xi64> loc(#loc48)
+    %x1_7 = arith.remsi %x1, %x1_6 : tensor<128xi64> loc(#loc48)
+    %x2 = tt.splat %ks2 : i64 -> tensor<128xi64> loc(#loc49)
+    %x2_8 = arith.divsi %x0, %x2 : tensor<128xi64> loc(#loc49)
+    %tmp0 = arith.cmpi sle, %ks1, %c1_i64 : i64 loc(#loc50)
+    %tmp0_9 = arith.cmpi sgt, %ks1, %c1_i64 : i64 loc(#loc51)
+    %tmp0_10 = arith.extui %tmp0_9 : i1 to i64 loc(#loc52)
+    %tmp0_11 = arith.muli %ks1, %tmp0_10 : i64 loc(#loc52)
+    %tmp0_12 = arith.extui %tmp0 : i1 to i64 loc(#loc68)
+    %tmp0_13 = arith.addi %tmp0_12, %tmp0_11 : i64 loc(#loc53)
+    %tmp0_14 = tt.splat %tmp0_13 : i64 -> tensor<128xi64> loc(#loc55)
+    %tmp0_15 = arith.muli %x0_5, %tmp0_14 : tensor<128xi64> loc(#loc55)
+    %tmp0_16 = arith.addi %x1_7, %tmp0_15 : tensor<128xi64> loc(#loc56)
+    %tmp0_17 = arith.cmpi sle, %ks0, %c1_i64 : i64 loc(#loc57)
+    %tmp0_18 = arith.cmpi sgt, %ks0, %c1_i64 : i64 loc(#loc58)
+    %tmp0_19 = arith.extui %tmp0_18 : i1 to i64 loc(#loc59)
+    %tmp0_20 = arith.muli %ks0, %tmp0_19 : i64 loc(#loc59)
+    %tmp0_21 = arith.extui %tmp0_17 : i1 to i64 loc(#loc69)
+    %tmp0_22 = arith.addi %tmp0_21, %tmp0_20 : i64 loc(#loc60)
+    %tmp0_23 = tt.splat %tmp0_22 : i64 -> tensor<128xi64> loc(#loc62)
+    %tmp0_24 = arith.muli %x2_8, %tmp0_23 : tensor<128xi64> loc(#loc62)
+    %tmp0_25 = arith.muli %tmp0_24, %tmp0_14 : tensor<128xi64> loc(#loc63)
+    %tmp0_26 = arith.addi %tmp0_16, %tmp0_25 : tensor<128xi64> loc(#loc64)
+    %tmp0_27 = tt.splat %in_ptr0 : !tt.ptr<i64> -> tensor<128x!tt.ptr<i64>> loc(#loc65)
+    %tmp0_28 = tt.addptr %tmp0_27, %tmp0_26 : tensor<128x!tt.ptr<i64>>, tensor<128xi64> loc(#loc65)
+    %tmp0_29 = tt.load %tmp0_28, %xmask_3 evictionPolicy = evict_last : tensor<128x!tt.ptr<i64>> loc(#loc66)
+    %tmp1 = arith.trunci %tmp0_29 : tensor<128xi64> to tensor<128xi32> loc(#loc67)
+    %0 = arith.muli %x1_7, %tmp0_23 : tensor<128xi64> loc(#loc29)
+    %1 = arith.addi %x0_5, %0 : tensor<128xi64> loc(#loc30)
+    %2 = arith.addi %1, %tmp0_25 : tensor<128xi64> loc(#loc31)
+    %3 = tt.splat %out_ptr0 : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>> loc(#loc32)
+    %4 = tt.addptr %3, %2 : tensor<128x!tt.ptr<i32>>, tensor<128xi64> loc(#loc32)
+    tt.store %4, %tmp1, %xmask_3 : tensor<128x!tt.ptr<i32>> loc(#loc33)
+    tt.return loc(#loc34)
+  } loc(#loc)
+} loc(#loc)
+#loc1 = loc(unknown)
+#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":19:28)
+#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":19:33)
+#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":20:36)
+#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":20:23)
+#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":21:21)
+#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":22:19)
+#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":23:21)
+#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":23:28)
+#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":24:19)
+#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:54)
+#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:80)
+#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:71)
+#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:62)
+#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:46)
+#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:39)
+#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:35)
+#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:106)
+#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:132)
+#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:123)
+#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:114)
+#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:98)
+#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:91)
+#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:138)
+#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:87)
+#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:30)
+#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":25:186)
+#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":26:19)
+#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:34)
+#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:30)
+#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:82)
+#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:25)
+#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:187)
+#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/35/c35huqp6ngzh67kt32kuxoqpghc32fstv4zogcouzabdxxwta3sl.py":27:4)
+#loc41 = loc("xoffset"(#loc2))
+#loc42 = loc("xoffset"(#loc3))
+#loc43 = loc("xindex"(#loc4))
+#loc44 = loc("xindex"(#loc5))
+#loc45 = loc("xmask"(#loc6))
+#loc46 = loc("x0"(#loc7))
+#loc47 = loc("x1"(#loc8))
+#loc48 = loc("x1"(#loc9))
+#loc49 = loc("x2"(#loc10))
+#loc50 = loc("tmp0"(#loc11))
+#loc51 = loc("tmp0"(#loc12))
+#loc52 = loc("tmp0"(#loc13))
+#loc53 = loc("tmp0"(#loc14))
+#loc54 = loc("tmp0"(#loc15))
+#loc55 = loc("tmp0"(#loc16))
+#loc56 = loc("tmp0"(#loc17))
+#loc57 = loc("tmp0"(#loc18))
+#loc58 = loc("tmp0"(#loc19))
+#loc59 = loc("tmp0"(#loc20))
+#loc60 = loc("tmp0"(#loc21))
+#loc61 = loc("tmp0"(#loc22))
+#loc62 = loc("tmp0"(#loc23))
+#loc63 = loc("tmp0"(#loc24))
+#loc64 = loc("tmp0"(#loc25))
+#loc65 = loc("tmp0"(#loc26))
+#loc66 = loc("tmp0"(#loc27))
+#loc67 = loc("tmp1"(#loc28))
+#loc68 = loc(fused[#loc53, #loc54])
+#loc69 = loc(fused[#loc60, #loc61])

SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/__grp__triton_red_fused_argmax_0.json

@@ -0,0 +1 @@
+{"child_paths": {"triton_red_fused_argmax_0.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.source", "triton_red_fused_argmax_0.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.ttir", "triton_red_fused_argmax_0.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.ttgir", "triton_red_fused_argmax_0.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.llir", "triton_red_fused_argmax_0.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.ptx", "triton_red_fused_argmax_0.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.cubin", "triton_red_fused_argmax_0.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/6/LEHCFAEILV2VVTHQWP55INWAP7WMDYDQAVN7JR5J546EKZSCPVWA/triton_red_fused_argmax_0.json"}}