Add files using upload-large-folder tool
Browse filesThis view is limited to 50 files because it contains too many changes. See raw diff
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/__grp__triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.cubin +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.llir +667 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ptx +1534 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.source +299 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttgir +232 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttir +231 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/__grp__triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.llir +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ptx +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.source +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttgir +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttir +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/__grp__triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.cubin +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.llir +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ptx +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.source +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttir +799 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/__grp__triton_poi_fused_new_zeros_1.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.cubin +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.llir +47 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ptx +207 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.source +41 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ttgir +35 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ttir +34 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/__grp__triton_red_fused__to_copy_mul_sum_0.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.cubin +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.llir +256 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ptx +688 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.source +338 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ttgir +176 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ttir +179 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/__grp__triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.cubin +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.llir +91 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ptx +373 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.source +206 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttgir +91 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttir +93 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/__grp__triton_poi_fused_new_zeros_1.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.cubin +0 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.json +1 -0
- SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.llir +47 -0
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/__grp__triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.source", "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttir", "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttgir", "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.llir", "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ptx", "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.cubin", "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.cubin
ADDED
|
Binary file (76 kB). View file
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"hash": "d4e9ec88be03aebb42041f53f2cbb9bb30afcd0e87a7a62c3ff7f5fe862eba44", "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_cat_index_mul_neg_slice_squeeze_unsqueeze_0"}
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.llir
ADDED
|
@@ -0,0 +1,667 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
; ModuleID = 'LLVMDialectModule'
|
| 2 |
+
source_filename = "LLVMDialectModule"
|
| 3 |
+
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"
|
| 4 |
+
|
| 5 |
+
@assertFunc_1 = internal constant [8 x i8] c"unknown\00"
|
| 6 |
+
@assertFile_1 = internal constant [114 x i8] c"/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py\00"
|
| 7 |
+
@assertMessage_1 = internal constant [38 x i8] c"index out of bounds: 0 <= tmp25 < ks4\00"
|
| 8 |
+
@assertFunc_0 = internal constant [8 x i8] c"unknown\00"
|
| 9 |
+
@assertFile_0 = internal constant [114 x i8] c"/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py\00"
|
| 10 |
+
@assertMessage_0 = internal constant [37 x i8] c"index out of bounds: 0 <= tmp5 < ks2\00"
|
| 11 |
+
|
| 12 |
+
; Function Attrs: noreturn
|
| 13 |
+
declare !dbg !5 void @__assertfail(ptr, ptr, i32, ptr, i64) local_unnamed_addr #0
|
| 14 |
+
|
| 15 |
+
define ptx_kernel void @triton_poi_fused_add_cat_index_mul_neg_slice_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, i64 %9, i32 %10, ptr addrspace(1) readnone captures(none) %11, ptr addrspace(1) readnone captures(none) %12) local_unnamed_addr #1 !dbg !9 {
|
| 16 |
+
%14 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !10
|
| 17 |
+
%15 = shl i32 %14, 10, !dbg !11
|
| 18 |
+
%16 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !12
|
| 19 |
+
%17 = shl nuw nsw i32 %16, 3, !dbg !12
|
| 20 |
+
%18 = and i32 %17, 1016, !dbg !12
|
| 21 |
+
%19 = or disjoint i32 %18, %15, !dbg !13
|
| 22 |
+
%20 = or disjoint i32 %19, 1, !dbg !13
|
| 23 |
+
%21 = or disjoint i32 %19, 2, !dbg !13
|
| 24 |
+
%22 = or disjoint i32 %19, 3, !dbg !13
|
| 25 |
+
%23 = or disjoint i32 %19, 4, !dbg !13
|
| 26 |
+
%24 = or disjoint i32 %19, 5, !dbg !13
|
| 27 |
+
%25 = or disjoint i32 %19, 6, !dbg !13
|
| 28 |
+
%26 = or disjoint i32 %19, 7, !dbg !13
|
| 29 |
+
%27 = insertelement <8 x i32> poison, i32 %26, i64 0, !dbg !14
|
| 30 |
+
%28 = insertelement <8 x i32> %27, i32 %25, i64 1, !dbg !14
|
| 31 |
+
%29 = insertelement <8 x i32> %28, i32 %24, i64 2, !dbg !14
|
| 32 |
+
%30 = insertelement <8 x i32> %29, i32 %23, i64 3, !dbg !14
|
| 33 |
+
%31 = insertelement <8 x i32> %30, i32 %22, i64 4, !dbg !14
|
| 34 |
+
%32 = insertelement <8 x i32> %31, i32 %21, i64 5, !dbg !14
|
| 35 |
+
%33 = insertelement <8 x i32> %32, i32 %20, i64 6, !dbg !14
|
| 36 |
+
%34 = insertelement <8 x i32> %33, i32 %19, i64 7, !dbg !14
|
| 37 |
+
%35 = sext <8 x i32> %34 to <8 x i64>, !dbg !14
|
| 38 |
+
%36 = extractelement <8 x i64> %35, i64 7, !dbg !15
|
| 39 |
+
%37 = sdiv i64 %36, %5, !dbg !14
|
| 40 |
+
%38 = extractelement <8 x i64> %35, i64 6, !dbg !15
|
| 41 |
+
%39 = sdiv i64 %38, %5, !dbg !14
|
| 42 |
+
%40 = extractelement <8 x i64> %35, i64 5, !dbg !15
|
| 43 |
+
%41 = sdiv i64 %40, %5, !dbg !14
|
| 44 |
+
%42 = extractelement <8 x i64> %35, i64 4, !dbg !15
|
| 45 |
+
%43 = sdiv i64 %42, %5, !dbg !14
|
| 46 |
+
%44 = extractelement <8 x i64> %35, i64 3, !dbg !15
|
| 47 |
+
%45 = sdiv i64 %44, %5, !dbg !14
|
| 48 |
+
%46 = extractelement <8 x i64> %35, i64 2, !dbg !15
|
| 49 |
+
%47 = sdiv i64 %46, %5, !dbg !14
|
| 50 |
+
%48 = extractelement <8 x i64> %35, i64 1, !dbg !15
|
| 51 |
+
%49 = sdiv i64 %48, %5, !dbg !14
|
| 52 |
+
%50 = extractelement <8 x i64> %35, i64 0, !dbg !15
|
| 53 |
+
%51 = sdiv i64 %50, %5, !dbg !14
|
| 54 |
+
%52 = srem i64 %37, %6, !dbg !16
|
| 55 |
+
%53 = srem i64 %39, %6, !dbg !16
|
| 56 |
+
%54 = srem i64 %41, %6, !dbg !16
|
| 57 |
+
%55 = srem i64 %43, %6, !dbg !16
|
| 58 |
+
%56 = srem i64 %45, %6, !dbg !16
|
| 59 |
+
%57 = srem i64 %47, %6, !dbg !16
|
| 60 |
+
%58 = srem i64 %49, %6, !dbg !16
|
| 61 |
+
%59 = srem i64 %51, %6, !dbg !16
|
| 62 |
+
%60 = getelementptr bfloat, ptr addrspace(1) %0, i64 %36, !dbg !15
|
| 63 |
+
%61 = getelementptr bfloat, ptr addrspace(1) %0, i64 %38, !dbg !15
|
| 64 |
+
%62 = getelementptr bfloat, ptr addrspace(1) %0, i64 %40, !dbg !15
|
| 65 |
+
%63 = getelementptr bfloat, ptr addrspace(1) %0, i64 %42, !dbg !15
|
| 66 |
+
%64 = getelementptr bfloat, ptr addrspace(1) %0, i64 %44, !dbg !15
|
| 67 |
+
%65 = getelementptr bfloat, ptr addrspace(1) %0, i64 %46, !dbg !15
|
| 68 |
+
%66 = getelementptr bfloat, ptr addrspace(1) %0, i64 %48, !dbg !15
|
| 69 |
+
%67 = getelementptr bfloat, ptr addrspace(1) %0, i64 %50, !dbg !15
|
| 70 |
+
%68 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 71 |
+
%69 = getelementptr i64, ptr addrspace(1) %1, i64 %52, !dbg !18
|
| 72 |
+
%70 = getelementptr i64, ptr addrspace(1) %1, i64 %53, !dbg !18
|
| 73 |
+
%71 = getelementptr i64, ptr addrspace(1) %1, i64 %54, !dbg !18
|
| 74 |
+
%72 = getelementptr i64, ptr addrspace(1) %1, i64 %55, !dbg !18
|
| 75 |
+
%73 = getelementptr i64, ptr addrspace(1) %1, i64 %56, !dbg !18
|
| 76 |
+
%74 = getelementptr i64, ptr addrspace(1) %1, i64 %57, !dbg !18
|
| 77 |
+
%75 = getelementptr i64, ptr addrspace(1) %1, i64 %58, !dbg !18
|
| 78 |
+
%76 = getelementptr i64, ptr addrspace(1) %1, i64 %59, !dbg !18
|
| 79 |
+
%77 = insertelement <8 x i32> poison, i32 %19, i64 0, !dbg !19
|
| 80 |
+
%78 = insertelement <8 x i32> %77, i32 %20, i64 1, !dbg !19
|
| 81 |
+
%79 = insertelement <8 x i32> %78, i32 %21, i64 2, !dbg !19
|
| 82 |
+
%80 = insertelement <8 x i32> %79, i32 %22, i64 3, !dbg !19
|
| 83 |
+
%81 = insertelement <8 x i32> %80, i32 %23, i64 4, !dbg !19
|
| 84 |
+
%82 = insertelement <8 x i32> %81, i32 %24, i64 5, !dbg !19
|
| 85 |
+
%83 = insertelement <8 x i32> %82, i32 %25, i64 6, !dbg !19
|
| 86 |
+
%84 = insertelement <8 x i32> %83, i32 %26, i64 7, !dbg !19
|
| 87 |
+
%85 = insertelement <8 x i32> poison, i32 %10, i64 0, !dbg !19
|
| 88 |
+
%86 = shufflevector <8 x i32> %85, <8 x i32> poison, <8 x i32> zeroinitializer, !dbg !19
|
| 89 |
+
%87 = icmp slt <8 x i32> %84, %86, !dbg !19
|
| 90 |
+
%88 = extractelement <8 x i1> %87, i64 0, !dbg !17
|
| 91 |
+
%89 = 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) %60, i64 %68, i1 %88) #4, !dbg !17
|
| 92 |
+
%90 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 93 |
+
%91 = extractelement <8 x i1> %87, i64 1, !dbg !17
|
| 94 |
+
%92 = 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) %61, i64 %90, i1 %91) #4, !dbg !17
|
| 95 |
+
%93 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 96 |
+
%94 = extractelement <8 x i1> %87, i64 2, !dbg !17
|
| 97 |
+
%95 = 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) %62, i64 %93, i1 %94) #4, !dbg !17
|
| 98 |
+
%96 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 99 |
+
%97 = extractelement <8 x i1> %87, i64 3, !dbg !17
|
| 100 |
+
%98 = 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) %63, i64 %96, i1 %97) #4, !dbg !17
|
| 101 |
+
%99 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 102 |
+
%100 = extractelement <8 x i1> %87, i64 4, !dbg !17
|
| 103 |
+
%101 = 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) %64, i64 %99, i1 %100) #4, !dbg !17
|
| 104 |
+
%102 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 105 |
+
%103 = extractelement <8 x i1> %87, i64 5, !dbg !17
|
| 106 |
+
%104 = 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 %102, i1 %103) #4, !dbg !17
|
| 107 |
+
%105 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 108 |
+
%106 = extractelement <8 x i1> %87, i64 6, !dbg !17
|
| 109 |
+
%107 = 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 %105, i1 %106) #4, !dbg !17
|
| 110 |
+
%108 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !17
|
| 111 |
+
%109 = extractelement <8 x i1> %87, i64 7, !dbg !17
|
| 112 |
+
%110 = 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 %108, i1 %109) #4, !dbg !17
|
| 113 |
+
%111 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 114 |
+
%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) %69, i64 %111, i1 %88) #4, !dbg !20
|
| 115 |
+
%113 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 116 |
+
%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) %70, i64 %113, i1 %91) #4, !dbg !20
|
| 117 |
+
%115 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 118 |
+
%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) %71, i64 %115, i1 %94) #4, !dbg !20
|
| 119 |
+
%117 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 120 |
+
%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) %72, i64 %117, i1 %97) #4, !dbg !20
|
| 121 |
+
%119 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 122 |
+
%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) %73, i64 %119, i1 %100) #4, !dbg !20
|
| 123 |
+
%121 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 124 |
+
%122 = 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) %74, i64 %121, i1 %103) #4, !dbg !20
|
| 125 |
+
%123 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 126 |
+
%124 = 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) %75, i64 %123, i1 %106) #4, !dbg !20
|
| 127 |
+
%125 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 128 |
+
%126 = 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) %76, i64 %125, i1 %109) #4, !dbg !20
|
| 129 |
+
%127 = insertelement <8 x i64> poison, i64 %112, i64 0, !dbg !21
|
| 130 |
+
%128 = insertelement <8 x i64> %127, i64 %114, i64 1, !dbg !21
|
| 131 |
+
%129 = insertelement <8 x i64> %128, i64 %116, i64 2, !dbg !21
|
| 132 |
+
%130 = insertelement <8 x i64> %129, i64 %118, i64 3, !dbg !21
|
| 133 |
+
%131 = insertelement <8 x i64> %130, i64 %120, i64 4, !dbg !21
|
| 134 |
+
%132 = insertelement <8 x i64> %131, i64 %122, i64 5, !dbg !21
|
| 135 |
+
%133 = insertelement <8 x i64> %132, i64 %124, i64 6, !dbg !21
|
| 136 |
+
%134 = insertelement <8 x i64> %133, i64 %126, i64 7, !dbg !21
|
| 137 |
+
%135 = icmp slt <8 x i64> %134, zeroinitializer, !dbg !21
|
| 138 |
+
%136 = insertelement <8 x i64> poison, i64 %7, i64 0, !dbg !22
|
| 139 |
+
%137 = shufflevector <8 x i64> %136, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !22
|
| 140 |
+
%138 = select <8 x i1> %135, <8 x i64> %137, <8 x i64> zeroinitializer, !dbg !22
|
| 141 |
+
%139 = add <8 x i64> %138, %134, !dbg !22
|
| 142 |
+
%140 = icmp slt <8 x i64> %139, zeroinitializer, !dbg !23
|
| 143 |
+
%141 = icmp sge <8 x i64> %139, %137, !dbg !24
|
| 144 |
+
%142 = or <8 x i1> %140, %141, !dbg !25
|
| 145 |
+
%143 = and <8 x i1> %87, %142, !dbg !26
|
| 146 |
+
%144 = bitcast <8 x i1> %143 to i8, !dbg !27
|
| 147 |
+
%.not = icmp eq i8 %144, 0, !dbg !27
|
| 148 |
+
br i1 %.not, label %146, label %145, !dbg !27
|
| 149 |
+
|
| 150 |
+
145: ; preds = %13
|
| 151 |
+
tail call void @__assertfail(ptr nonnull @assertMessage_0, ptr nonnull @assertFile_0, i32 32, ptr nonnull @assertFunc_0, i64 1), !dbg !27
|
| 152 |
+
unreachable, !dbg !27
|
| 153 |
+
|
| 154 |
+
146: ; preds = %13
|
| 155 |
+
%147 = insertelement <8 x i64> poison, i64 %8, i64 0, !dbg !28
|
| 156 |
+
%148 = shufflevector <8 x i64> %147, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !28
|
| 157 |
+
%149 = srem <8 x i64> %35, %148, !dbg !28
|
| 158 |
+
tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !27
|
| 159 |
+
%150 = extractelement <8 x i64> %139, i64 0, !dbg !29
|
| 160 |
+
%151 = mul i64 %150, %8, !dbg !29
|
| 161 |
+
%152 = extractelement <8 x i64> %139, i64 1, !dbg !29
|
| 162 |
+
%153 = mul i64 %152, %8, !dbg !29
|
| 163 |
+
%154 = extractelement <8 x i64> %139, i64 2, !dbg !29
|
| 164 |
+
%155 = mul i64 %154, %8, !dbg !29
|
| 165 |
+
%156 = extractelement <8 x i64> %139, i64 3, !dbg !29
|
| 166 |
+
%157 = mul i64 %156, %8, !dbg !29
|
| 167 |
+
%158 = extractelement <8 x i64> %139, i64 4, !dbg !29
|
| 168 |
+
%159 = mul i64 %158, %8, !dbg !29
|
| 169 |
+
%160 = extractelement <8 x i64> %139, i64 5, !dbg !29
|
| 170 |
+
%161 = mul i64 %160, %8, !dbg !29
|
| 171 |
+
%162 = extractelement <8 x i64> %139, i64 6, !dbg !29
|
| 172 |
+
%163 = mul i64 %162, %8, !dbg !29
|
| 173 |
+
%164 = extractelement <8 x i64> %139, i64 7, !dbg !29
|
| 174 |
+
%165 = mul i64 %164, %8, !dbg !29
|
| 175 |
+
%166 = extractelement <8 x i64> %149, i64 7, !dbg !30
|
| 176 |
+
%167 = getelementptr bfloat, ptr addrspace(1) %2, i64 %166, !dbg !31
|
| 177 |
+
%168 = getelementptr bfloat, ptr addrspace(1) %167, i64 %151, !dbg !31
|
| 178 |
+
%169 = extractelement <8 x i64> %149, i64 6, !dbg !30
|
| 179 |
+
%170 = getelementptr bfloat, ptr addrspace(1) %2, i64 %169, !dbg !31
|
| 180 |
+
%171 = getelementptr bfloat, ptr addrspace(1) %170, i64 %153, !dbg !31
|
| 181 |
+
%172 = extractelement <8 x i64> %149, i64 5, !dbg !30
|
| 182 |
+
%173 = getelementptr bfloat, ptr addrspace(1) %2, i64 %172, !dbg !31
|
| 183 |
+
%174 = getelementptr bfloat, ptr addrspace(1) %173, i64 %155, !dbg !31
|
| 184 |
+
%175 = extractelement <8 x i64> %149, i64 4, !dbg !30
|
| 185 |
+
%176 = getelementptr bfloat, ptr addrspace(1) %2, i64 %175, !dbg !31
|
| 186 |
+
%177 = getelementptr bfloat, ptr addrspace(1) %176, i64 %157, !dbg !31
|
| 187 |
+
%178 = extractelement <8 x i64> %149, i64 3, !dbg !30
|
| 188 |
+
%179 = getelementptr bfloat, ptr addrspace(1) %2, i64 %178, !dbg !31
|
| 189 |
+
%180 = getelementptr bfloat, ptr addrspace(1) %179, i64 %159, !dbg !31
|
| 190 |
+
%181 = extractelement <8 x i64> %149, i64 2, !dbg !30
|
| 191 |
+
%182 = getelementptr bfloat, ptr addrspace(1) %2, i64 %181, !dbg !31
|
| 192 |
+
%183 = getelementptr bfloat, ptr addrspace(1) %182, i64 %161, !dbg !31
|
| 193 |
+
%184 = extractelement <8 x i64> %149, i64 1, !dbg !30
|
| 194 |
+
%185 = getelementptr bfloat, ptr addrspace(1) %2, i64 %184, !dbg !31
|
| 195 |
+
%186 = getelementptr bfloat, ptr addrspace(1) %185, i64 %163, !dbg !31
|
| 196 |
+
%187 = extractelement <8 x i64> %149, i64 0, !dbg !30
|
| 197 |
+
%188 = getelementptr bfloat, ptr addrspace(1) %2, i64 %187, !dbg !31
|
| 198 |
+
%189 = getelementptr bfloat, ptr addrspace(1) %188, i64 %165, !dbg !31
|
| 199 |
+
%190 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 200 |
+
%191 = 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) %168, i64 %190, i1 %88) #4, !dbg !32
|
| 201 |
+
%192 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 202 |
+
%193 = 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) %171, i64 %192, i1 %91) #4, !dbg !32
|
| 203 |
+
%194 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 204 |
+
%195 = 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) %174, i64 %194, i1 %94) #4, !dbg !32
|
| 205 |
+
%196 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 206 |
+
%197 = 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) %177, i64 %196, i1 %97) #4, !dbg !32
|
| 207 |
+
%198 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 208 |
+
%199 = 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) %180, i64 %198, i1 %100) #4, !dbg !32
|
| 209 |
+
%200 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 210 |
+
%201 = 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) %183, i64 %200, i1 %103) #4, !dbg !32
|
| 211 |
+
%202 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 212 |
+
%203 = 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) %186, i64 %202, i1 %106) #4, !dbg !32
|
| 213 |
+
%204 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !32
|
| 214 |
+
%205 = 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) %189, i64 %204, i1 %109) #4, !dbg !32
|
| 215 |
+
%206 = sdiv i64 %8, 2, !dbg !33
|
| 216 |
+
%207 = sub i64 %8, %206, !dbg !34
|
| 217 |
+
%208 = icmp slt i64 %166, %207, !dbg !35
|
| 218 |
+
%209 = icmp slt i64 %169, %207, !dbg !35
|
| 219 |
+
%210 = icmp slt i64 %172, %207, !dbg !35
|
| 220 |
+
%211 = icmp slt i64 %175, %207, !dbg !35
|
| 221 |
+
%212 = icmp slt i64 %178, %207, !dbg !35
|
| 222 |
+
%213 = icmp slt i64 %181, %207, !dbg !35
|
| 223 |
+
%214 = icmp slt i64 %184, %207, !dbg !35
|
| 224 |
+
%215 = icmp slt i64 %187, %207, !dbg !35
|
| 225 |
+
%216 = sub nsw i64 %36, %166, !dbg !30
|
| 226 |
+
%217 = sub nsw i64 %38, %169, !dbg !30
|
| 227 |
+
%218 = sub nsw i64 %40, %172, !dbg !30
|
| 228 |
+
%219 = sub nsw i64 %42, %175, !dbg !30
|
| 229 |
+
%220 = sub nsw i64 %44, %178, !dbg !30
|
| 230 |
+
%221 = sub nsw i64 %46, %181, !dbg !30
|
| 231 |
+
%222 = sub nsw i64 %48, %184, !dbg !30
|
| 232 |
+
%223 = sub nsw i64 %50, %187, !dbg !30
|
| 233 |
+
%224 = getelementptr bfloat, ptr addrspace(1) %0, i64 %216, !dbg !36
|
| 234 |
+
%225 = getelementptr bfloat, ptr addrspace(1) %224, i64 %206, !dbg !36
|
| 235 |
+
%226 = getelementptr bfloat, ptr addrspace(1) %225, i64 %166, !dbg !36
|
| 236 |
+
%227 = getelementptr bfloat, ptr addrspace(1) %0, i64 %217, !dbg !36
|
| 237 |
+
%228 = getelementptr bfloat, ptr addrspace(1) %227, i64 %206, !dbg !36
|
| 238 |
+
%229 = getelementptr bfloat, ptr addrspace(1) %228, i64 %169, !dbg !36
|
| 239 |
+
%230 = getelementptr bfloat, ptr addrspace(1) %0, i64 %218, !dbg !36
|
| 240 |
+
%231 = getelementptr bfloat, ptr addrspace(1) %230, i64 %206, !dbg !36
|
| 241 |
+
%232 = getelementptr bfloat, ptr addrspace(1) %231, i64 %172, !dbg !36
|
| 242 |
+
%233 = getelementptr bfloat, ptr addrspace(1) %0, i64 %219, !dbg !36
|
| 243 |
+
%234 = getelementptr bfloat, ptr addrspace(1) %233, i64 %206, !dbg !36
|
| 244 |
+
%235 = getelementptr bfloat, ptr addrspace(1) %234, i64 %175, !dbg !36
|
| 245 |
+
%236 = getelementptr bfloat, ptr addrspace(1) %0, i64 %220, !dbg !36
|
| 246 |
+
%237 = getelementptr bfloat, ptr addrspace(1) %236, i64 %206, !dbg !36
|
| 247 |
+
%238 = getelementptr bfloat, ptr addrspace(1) %237, i64 %178, !dbg !36
|
| 248 |
+
%239 = getelementptr bfloat, ptr addrspace(1) %0, i64 %221, !dbg !36
|
| 249 |
+
%240 = getelementptr bfloat, ptr addrspace(1) %239, i64 %206, !dbg !36
|
| 250 |
+
%241 = getelementptr bfloat, ptr addrspace(1) %240, i64 %181, !dbg !36
|
| 251 |
+
%242 = getelementptr bfloat, ptr addrspace(1) %0, i64 %222, !dbg !36
|
| 252 |
+
%243 = getelementptr bfloat, ptr addrspace(1) %242, i64 %206, !dbg !36
|
| 253 |
+
%244 = getelementptr bfloat, ptr addrspace(1) %243, i64 %184, !dbg !36
|
| 254 |
+
%245 = getelementptr bfloat, ptr addrspace(1) %0, i64 %223, !dbg !36
|
| 255 |
+
%246 = getelementptr bfloat, ptr addrspace(1) %245, i64 %206, !dbg !36
|
| 256 |
+
%247 = getelementptr bfloat, ptr addrspace(1) %246, i64 %187, !dbg !36
|
| 257 |
+
%248 = and i1 %88, %208, !dbg !37
|
| 258 |
+
%249 = and i1 %91, %209, !dbg !37
|
| 259 |
+
%250 = and i1 %94, %210, !dbg !37
|
| 260 |
+
%251 = and i1 %97, %211, !dbg !37
|
| 261 |
+
%252 = and i1 %100, %212, !dbg !37
|
| 262 |
+
%253 = and i1 %103, %213, !dbg !37
|
| 263 |
+
%254 = and i1 %106, %214, !dbg !37
|
| 264 |
+
%255 = and i1 %109, %215, !dbg !37
|
| 265 |
+
%256 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 266 |
+
%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) %226, i64 %256, i1 %248) #4, !dbg !38
|
| 267 |
+
%258 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 268 |
+
%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) %229, i64 %258, i1 %249) #4, !dbg !38
|
| 269 |
+
%260 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 270 |
+
%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) %232, i64 %260, i1 %250) #4, !dbg !38
|
| 271 |
+
%262 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 272 |
+
%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) %235, i64 %262, i1 %251) #4, !dbg !38
|
| 273 |
+
%264 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 274 |
+
%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) %238, i64 %264, i1 %252) #4, !dbg !38
|
| 275 |
+
%266 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 276 |
+
%267 = 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 %266, i1 %253) #4, !dbg !38
|
| 277 |
+
%268 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 278 |
+
%269 = 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) %244, i64 %268, i1 %254) #4, !dbg !38
|
| 279 |
+
%270 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !38
|
| 280 |
+
%271 = 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 %270, i1 %255) #4, !dbg !38
|
| 281 |
+
%272 = insertelement <8 x i64> poison, i64 %207, i64 0, !dbg !39
|
| 282 |
+
%273 = shufflevector <8 x i64> %272, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !39
|
| 283 |
+
%274 = icmp sge <8 x i64> %149, %273, !dbg !39
|
| 284 |
+
%275 = sub i64 %166, %8, !dbg !40
|
| 285 |
+
%276 = sub i64 %169, %8, !dbg !40
|
| 286 |
+
%277 = sub i64 %172, %8, !dbg !40
|
| 287 |
+
%278 = sub i64 %175, %8, !dbg !40
|
| 288 |
+
%279 = sub i64 %178, %8, !dbg !40
|
| 289 |
+
%280 = sub i64 %181, %8, !dbg !40
|
| 290 |
+
%281 = sub i64 %184, %8, !dbg !40
|
| 291 |
+
%282 = sub i64 %187, %8, !dbg !40
|
| 292 |
+
%283 = getelementptr bfloat, ptr addrspace(1) %224, i64 %275, !dbg !41
|
| 293 |
+
%284 = getelementptr bfloat, ptr addrspace(1) %283, i64 %206, !dbg !41
|
| 294 |
+
%285 = getelementptr bfloat, ptr addrspace(1) %227, i64 %276, !dbg !41
|
| 295 |
+
%286 = getelementptr bfloat, ptr addrspace(1) %285, i64 %206, !dbg !41
|
| 296 |
+
%287 = getelementptr bfloat, ptr addrspace(1) %230, i64 %277, !dbg !41
|
| 297 |
+
%288 = getelementptr bfloat, ptr addrspace(1) %287, i64 %206, !dbg !41
|
| 298 |
+
%289 = getelementptr bfloat, ptr addrspace(1) %233, i64 %278, !dbg !41
|
| 299 |
+
%290 = getelementptr bfloat, ptr addrspace(1) %289, i64 %206, !dbg !41
|
| 300 |
+
%291 = getelementptr bfloat, ptr addrspace(1) %236, i64 %279, !dbg !41
|
| 301 |
+
%292 = getelementptr bfloat, ptr addrspace(1) %291, i64 %206, !dbg !41
|
| 302 |
+
%293 = getelementptr bfloat, ptr addrspace(1) %239, i64 %280, !dbg !41
|
| 303 |
+
%294 = getelementptr bfloat, ptr addrspace(1) %293, i64 %206, !dbg !41
|
| 304 |
+
%295 = getelementptr bfloat, ptr addrspace(1) %242, i64 %281, !dbg !41
|
| 305 |
+
%296 = getelementptr bfloat, ptr addrspace(1) %295, i64 %206, !dbg !41
|
| 306 |
+
%297 = getelementptr bfloat, ptr addrspace(1) %245, i64 %282, !dbg !41
|
| 307 |
+
%298 = getelementptr bfloat, ptr addrspace(1) %297, i64 %206, !dbg !41
|
| 308 |
+
%299 = extractelement <8 x i1> %274, i64 7, !dbg !42
|
| 309 |
+
%300 = and i1 %88, %299, !dbg !42
|
| 310 |
+
%301 = extractelement <8 x i1> %274, i64 6, !dbg !42
|
| 311 |
+
%302 = and i1 %91, %301, !dbg !42
|
| 312 |
+
%303 = extractelement <8 x i1> %274, i64 5, !dbg !42
|
| 313 |
+
%304 = and i1 %94, %303, !dbg !42
|
| 314 |
+
%305 = extractelement <8 x i1> %274, i64 4, !dbg !42
|
| 315 |
+
%306 = and i1 %97, %305, !dbg !42
|
| 316 |
+
%307 = extractelement <8 x i1> %274, i64 3, !dbg !42
|
| 317 |
+
%308 = and i1 %100, %307, !dbg !42
|
| 318 |
+
%309 = extractelement <8 x i1> %274, i64 2, !dbg !42
|
| 319 |
+
%310 = and i1 %103, %309, !dbg !42
|
| 320 |
+
%311 = extractelement <8 x i1> %274, i64 1, !dbg !42
|
| 321 |
+
%312 = and i1 %106, %311, !dbg !42
|
| 322 |
+
%313 = extractelement <8 x i1> %274, i64 0, !dbg !42
|
| 323 |
+
%314 = and i1 %109, %313, !dbg !42
|
| 324 |
+
%315 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 325 |
+
%316 = 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 %315, i1 %300) #4, !dbg !43
|
| 326 |
+
%317 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 327 |
+
%318 = 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 %317, i1 %302) #4, !dbg !43
|
| 328 |
+
%319 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 329 |
+
%320 = 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 %319, i1 %304) #4, !dbg !43
|
| 330 |
+
%321 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 331 |
+
%322 = 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 %321, i1 %306) #4, !dbg !43
|
| 332 |
+
%323 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 333 |
+
%324 = 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) %292, i64 %323, i1 %308) #4, !dbg !43
|
| 334 |
+
%325 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 335 |
+
%326 = 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) %294, i64 %325, i1 %310) #4, !dbg !43
|
| 336 |
+
%327 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 337 |
+
%328 = 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) %296, i64 %327, i1 %312) #4, !dbg !43
|
| 338 |
+
%329 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !43
|
| 339 |
+
%330 = 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) %298, i64 %329, i1 %314) #4, !dbg !43
|
| 340 |
+
%331 = insertelement <8 x i64> poison, i64 %9, i64 0, !dbg !44
|
| 341 |
+
%332 = shufflevector <8 x i64> %331, <8 x i64> poison, <8 x i32> zeroinitializer, !dbg !44
|
| 342 |
+
%333 = select <8 x i1> %135, <8 x i64> %332, <8 x i64> zeroinitializer, !dbg !44
|
| 343 |
+
%334 = add <8 x i64> %333, %134, !dbg !44
|
| 344 |
+
%335 = icmp slt <8 x i64> %334, zeroinitializer, !dbg !45
|
| 345 |
+
%336 = icmp sge <8 x i64> %334, %332, !dbg !46
|
| 346 |
+
%337 = or <8 x i1> %335, %336, !dbg !47
|
| 347 |
+
%338 = and <8 x i1> %87, %337, !dbg !48
|
| 348 |
+
%339 = bitcast <8 x i1> %338 to i8, !dbg !49
|
| 349 |
+
%.not87 = icmp eq i8 %339, 0, !dbg !49
|
| 350 |
+
br i1 %.not87, label %341, label %340, !dbg !49
|
| 351 |
+
|
| 352 |
+
340: ; preds = %146
|
| 353 |
+
tail call void @__assertfail(ptr nonnull @assertMessage_1, ptr nonnull @assertFile_1, i32 52, ptr nonnull @assertFunc_1, i64 1), !dbg !49
|
| 354 |
+
unreachable, !dbg !49
|
| 355 |
+
|
| 356 |
+
341: ; preds = %146
|
| 357 |
+
%342 = bitcast i16 %271 to bfloat, !dbg !38
|
| 358 |
+
%343 = fpext bfloat %342 to float, !dbg !50
|
| 359 |
+
%344 = fsub float 0.000000e+00, %343, !dbg !51
|
| 360 |
+
%345 = bitcast i16 %330 to bfloat, !dbg !43
|
| 361 |
+
%346 = fpext bfloat %345 to float, !dbg !52
|
| 362 |
+
%347 = select i1 %215, float %344, float %346, !dbg !53
|
| 363 |
+
%348 = bitcast i16 %269 to bfloat, !dbg !38
|
| 364 |
+
%349 = fpext bfloat %348 to float, !dbg !50
|
| 365 |
+
%350 = fsub float 0.000000e+00, %349, !dbg !51
|
| 366 |
+
%351 = bitcast i16 %328 to bfloat, !dbg !43
|
| 367 |
+
%352 = fpext bfloat %351 to float, !dbg !52
|
| 368 |
+
%353 = select i1 %214, float %350, float %352, !dbg !53
|
| 369 |
+
%354 = bitcast i16 %267 to bfloat, !dbg !38
|
| 370 |
+
%355 = fpext bfloat %354 to float, !dbg !50
|
| 371 |
+
%356 = fsub float 0.000000e+00, %355, !dbg !51
|
| 372 |
+
%357 = bitcast i16 %326 to bfloat, !dbg !43
|
| 373 |
+
%358 = fpext bfloat %357 to float, !dbg !52
|
| 374 |
+
%359 = select i1 %213, float %356, float %358, !dbg !53
|
| 375 |
+
%360 = bitcast i16 %265 to bfloat, !dbg !38
|
| 376 |
+
%361 = fpext bfloat %360 to float, !dbg !50
|
| 377 |
+
%362 = fsub float 0.000000e+00, %361, !dbg !51
|
| 378 |
+
%363 = bitcast i16 %324 to bfloat, !dbg !43
|
| 379 |
+
%364 = fpext bfloat %363 to float, !dbg !52
|
| 380 |
+
%365 = select i1 %212, float %362, float %364, !dbg !53
|
| 381 |
+
%366 = bitcast i16 %263 to bfloat, !dbg !38
|
| 382 |
+
%367 = fpext bfloat %366 to float, !dbg !50
|
| 383 |
+
%368 = fsub float 0.000000e+00, %367, !dbg !51
|
| 384 |
+
%369 = bitcast i16 %322 to bfloat, !dbg !43
|
| 385 |
+
%370 = fpext bfloat %369 to float, !dbg !52
|
| 386 |
+
%371 = select i1 %211, float %368, float %370, !dbg !53
|
| 387 |
+
%372 = bitcast i16 %261 to bfloat, !dbg !38
|
| 388 |
+
%373 = fpext bfloat %372 to float, !dbg !50
|
| 389 |
+
%374 = fsub float 0.000000e+00, %373, !dbg !51
|
| 390 |
+
%375 = bitcast i16 %320 to bfloat, !dbg !43
|
| 391 |
+
%376 = fpext bfloat %375 to float, !dbg !52
|
| 392 |
+
%377 = select i1 %210, float %374, float %376, !dbg !53
|
| 393 |
+
%378 = bitcast i16 %259 to bfloat, !dbg !38
|
| 394 |
+
%379 = fpext bfloat %378 to float, !dbg !50
|
| 395 |
+
%380 = fsub float 0.000000e+00, %379, !dbg !51
|
| 396 |
+
%381 = bitcast i16 %318 to bfloat, !dbg !43
|
| 397 |
+
%382 = fpext bfloat %381 to float, !dbg !52
|
| 398 |
+
%383 = select i1 %209, float %380, float %382, !dbg !53
|
| 399 |
+
%384 = bitcast i16 %257 to bfloat, !dbg !38
|
| 400 |
+
%385 = fpext bfloat %384 to float, !dbg !50
|
| 401 |
+
%386 = fsub float 0.000000e+00, %385, !dbg !51
|
| 402 |
+
%387 = bitcast i16 %316 to bfloat, !dbg !43
|
| 403 |
+
%388 = fpext bfloat %387 to float, !dbg !52
|
| 404 |
+
%389 = select i1 %208, float %386, float %388, !dbg !53
|
| 405 |
+
%390 = bitcast i16 %110 to bfloat, !dbg !17
|
| 406 |
+
%391 = fpext bfloat %390 to float, !dbg !54
|
| 407 |
+
%392 = bitcast i16 %107 to bfloat, !dbg !17
|
| 408 |
+
%393 = fpext bfloat %392 to float, !dbg !54
|
| 409 |
+
%394 = bitcast i16 %104 to bfloat, !dbg !17
|
| 410 |
+
%395 = fpext bfloat %394 to float, !dbg !54
|
| 411 |
+
%396 = bitcast i16 %101 to bfloat, !dbg !17
|
| 412 |
+
%397 = fpext bfloat %396 to float, !dbg !54
|
| 413 |
+
%398 = bitcast i16 %98 to bfloat, !dbg !17
|
| 414 |
+
%399 = fpext bfloat %398 to float, !dbg !54
|
| 415 |
+
%400 = bitcast i16 %95 to bfloat, !dbg !17
|
| 416 |
+
%401 = fpext bfloat %400 to float, !dbg !54
|
| 417 |
+
%402 = bitcast i16 %92 to bfloat, !dbg !17
|
| 418 |
+
%403 = fpext bfloat %402 to float, !dbg !54
|
| 419 |
+
%404 = bitcast i16 %89 to bfloat, !dbg !17
|
| 420 |
+
%405 = fpext bfloat %404 to float, !dbg !54
|
| 421 |
+
tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !49
|
| 422 |
+
%406 = extractelement <8 x i64> %334, i64 0, !dbg !55
|
| 423 |
+
%407 = mul i64 %406, %8, !dbg !55
|
| 424 |
+
%408 = extractelement <8 x i64> %334, i64 1, !dbg !55
|
| 425 |
+
%409 = mul i64 %408, %8, !dbg !55
|
| 426 |
+
%410 = extractelement <8 x i64> %334, i64 2, !dbg !55
|
| 427 |
+
%411 = mul i64 %410, %8, !dbg !55
|
| 428 |
+
%412 = extractelement <8 x i64> %334, i64 3, !dbg !55
|
| 429 |
+
%413 = mul i64 %412, %8, !dbg !55
|
| 430 |
+
%414 = extractelement <8 x i64> %334, i64 4, !dbg !55
|
| 431 |
+
%415 = mul i64 %414, %8, !dbg !55
|
| 432 |
+
%416 = extractelement <8 x i64> %334, i64 5, !dbg !55
|
| 433 |
+
%417 = mul i64 %416, %8, !dbg !55
|
| 434 |
+
%418 = extractelement <8 x i64> %334, i64 6, !dbg !55
|
| 435 |
+
%419 = mul i64 %418, %8, !dbg !55
|
| 436 |
+
%420 = extractelement <8 x i64> %334, i64 7, !dbg !55
|
| 437 |
+
%421 = mul i64 %420, %8, !dbg !55
|
| 438 |
+
%422 = getelementptr bfloat, ptr addrspace(1) %3, i64 %166, !dbg !56
|
| 439 |
+
%423 = getelementptr bfloat, ptr addrspace(1) %422, i64 %407, !dbg !56
|
| 440 |
+
%424 = getelementptr bfloat, ptr addrspace(1) %3, i64 %169, !dbg !56
|
| 441 |
+
%425 = getelementptr bfloat, ptr addrspace(1) %424, i64 %409, !dbg !56
|
| 442 |
+
%426 = getelementptr bfloat, ptr addrspace(1) %3, i64 %172, !dbg !56
|
| 443 |
+
%427 = getelementptr bfloat, ptr addrspace(1) %426, i64 %411, !dbg !56
|
| 444 |
+
%428 = getelementptr bfloat, ptr addrspace(1) %3, i64 %175, !dbg !56
|
| 445 |
+
%429 = getelementptr bfloat, ptr addrspace(1) %428, i64 %413, !dbg !56
|
| 446 |
+
%430 = getelementptr bfloat, ptr addrspace(1) %3, i64 %178, !dbg !56
|
| 447 |
+
%431 = getelementptr bfloat, ptr addrspace(1) %430, i64 %415, !dbg !56
|
| 448 |
+
%432 = getelementptr bfloat, ptr addrspace(1) %3, i64 %181, !dbg !56
|
| 449 |
+
%433 = getelementptr bfloat, ptr addrspace(1) %432, i64 %417, !dbg !56
|
| 450 |
+
%434 = getelementptr bfloat, ptr addrspace(1) %3, i64 %184, !dbg !56
|
| 451 |
+
%435 = getelementptr bfloat, ptr addrspace(1) %434, i64 %419, !dbg !56
|
| 452 |
+
%436 = getelementptr bfloat, ptr addrspace(1) %3, i64 %187, !dbg !56
|
| 453 |
+
%437 = getelementptr bfloat, ptr addrspace(1) %436, i64 %421, !dbg !56
|
| 454 |
+
%438 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 455 |
+
%439 = 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) %423, i64 %438, i1 %88) #4, !dbg !57
|
| 456 |
+
%440 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 457 |
+
%441 = 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) %425, i64 %440, i1 %91) #4, !dbg !57
|
| 458 |
+
%442 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 459 |
+
%443 = 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) %427, i64 %442, i1 %94) #4, !dbg !57
|
| 460 |
+
%444 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 461 |
+
%445 = 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) %429, i64 %444, i1 %97) #4, !dbg !57
|
| 462 |
+
%446 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 463 |
+
%447 = 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) %431, i64 %446, i1 %100) #4, !dbg !57
|
| 464 |
+
%448 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 465 |
+
%449 = 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) %433, i64 %448, i1 %103) #4, !dbg !57
|
| 466 |
+
%450 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 467 |
+
%451 = 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) %435, i64 %450, i1 %106) #4, !dbg !57
|
| 468 |
+
%452 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !57
|
| 469 |
+
%453 = 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) %437, i64 %452, i1 %109) #4, !dbg !57
|
| 470 |
+
%454 = insertelement <2 x i16> poison, i16 %191, i64 0, !dbg !32
|
| 471 |
+
%455 = insertelement <2 x i16> %454, i16 %439, i64 1, !dbg !32
|
| 472 |
+
%456 = bitcast <2 x i16> %455 to <2 x bfloat>, !dbg !32
|
| 473 |
+
%457 = fpext <2 x bfloat> %456 to <2 x float>, !dbg !58
|
| 474 |
+
%458 = insertelement <2 x float> poison, float %405, i64 0, !dbg !59
|
| 475 |
+
%459 = insertelement <2 x float> %458, float %389, i64 1, !dbg !59
|
| 476 |
+
%460 = fmul <2 x float> %459, %457, !dbg !59
|
| 477 |
+
%461 = insertelement <2 x i16> poison, i16 %193, i64 0, !dbg !32
|
| 478 |
+
%462 = insertelement <2 x i16> %461, i16 %441, i64 1, !dbg !32
|
| 479 |
+
%463 = bitcast <2 x i16> %462 to <2 x bfloat>, !dbg !32
|
| 480 |
+
%464 = fpext <2 x bfloat> %463 to <2 x float>, !dbg !58
|
| 481 |
+
%465 = insertelement <2 x float> poison, float %403, i64 0, !dbg !59
|
| 482 |
+
%466 = insertelement <2 x float> %465, float %383, i64 1, !dbg !59
|
| 483 |
+
%467 = fmul <2 x float> %466, %464, !dbg !59
|
| 484 |
+
%468 = insertelement <2 x i16> poison, i16 %195, i64 0, !dbg !32
|
| 485 |
+
%469 = insertelement <2 x i16> %468, i16 %443, i64 1, !dbg !32
|
| 486 |
+
%470 = bitcast <2 x i16> %469 to <2 x bfloat>, !dbg !32
|
| 487 |
+
%471 = fpext <2 x bfloat> %470 to <2 x float>, !dbg !58
|
| 488 |
+
%472 = insertelement <2 x float> poison, float %401, i64 0, !dbg !59
|
| 489 |
+
%473 = insertelement <2 x float> %472, float %377, i64 1, !dbg !59
|
| 490 |
+
%474 = fmul <2 x float> %473, %471, !dbg !59
|
| 491 |
+
%475 = insertelement <2 x i16> poison, i16 %197, i64 0, !dbg !32
|
| 492 |
+
%476 = insertelement <2 x i16> %475, i16 %445, i64 1, !dbg !32
|
| 493 |
+
%477 = bitcast <2 x i16> %476 to <2 x bfloat>, !dbg !32
|
| 494 |
+
%478 = fpext <2 x bfloat> %477 to <2 x float>, !dbg !58
|
| 495 |
+
%479 = insertelement <2 x float> poison, float %399, i64 0, !dbg !59
|
| 496 |
+
%480 = insertelement <2 x float> %479, float %371, i64 1, !dbg !59
|
| 497 |
+
%481 = fmul <2 x float> %480, %478, !dbg !59
|
| 498 |
+
%482 = insertelement <2 x i16> poison, i16 %199, i64 0, !dbg !32
|
| 499 |
+
%483 = insertelement <2 x i16> %482, i16 %447, i64 1, !dbg !32
|
| 500 |
+
%484 = bitcast <2 x i16> %483 to <2 x bfloat>, !dbg !32
|
| 501 |
+
%485 = fpext <2 x bfloat> %484 to <2 x float>, !dbg !58
|
| 502 |
+
%486 = insertelement <2 x float> poison, float %397, i64 0, !dbg !59
|
| 503 |
+
%487 = insertelement <2 x float> %486, float %365, i64 1, !dbg !59
|
| 504 |
+
%488 = fmul <2 x float> %487, %485, !dbg !59
|
| 505 |
+
%489 = insertelement <2 x i16> poison, i16 %201, i64 0, !dbg !32
|
| 506 |
+
%490 = insertelement <2 x i16> %489, i16 %449, i64 1, !dbg !32
|
| 507 |
+
%491 = bitcast <2 x i16> %490 to <2 x bfloat>, !dbg !32
|
| 508 |
+
%492 = fpext <2 x bfloat> %491 to <2 x float>, !dbg !58
|
| 509 |
+
%493 = insertelement <2 x float> poison, float %395, i64 0, !dbg !59
|
| 510 |
+
%494 = insertelement <2 x float> %493, float %359, i64 1, !dbg !59
|
| 511 |
+
%495 = fmul <2 x float> %494, %492, !dbg !59
|
| 512 |
+
%496 = insertelement <2 x i16> poison, i16 %203, i64 0, !dbg !32
|
| 513 |
+
%497 = insertelement <2 x i16> %496, i16 %451, i64 1, !dbg !32
|
| 514 |
+
%498 = bitcast <2 x i16> %497 to <2 x bfloat>, !dbg !32
|
| 515 |
+
%499 = fpext <2 x bfloat> %498 to <2 x float>, !dbg !58
|
| 516 |
+
%500 = insertelement <2 x float> poison, float %393, i64 0, !dbg !59
|
| 517 |
+
%501 = insertelement <2 x float> %500, float %353, i64 1, !dbg !59
|
| 518 |
+
%502 = fmul <2 x float> %501, %499, !dbg !59
|
| 519 |
+
%503 = insertelement <2 x i16> poison, i16 %205, i64 0, !dbg !32
|
| 520 |
+
%504 = insertelement <2 x i16> %503, i16 %453, i64 1, !dbg !32
|
| 521 |
+
%505 = bitcast <2 x i16> %504 to <2 x bfloat>, !dbg !32
|
| 522 |
+
%506 = fpext <2 x bfloat> %505 to <2 x float>, !dbg !58
|
| 523 |
+
%507 = insertelement <2 x float> poison, float %391, i64 0, !dbg !59
|
| 524 |
+
%508 = insertelement <2 x float> %507, float %347, i64 1, !dbg !59
|
| 525 |
+
%509 = fmul <2 x float> %508, %506, !dbg !59
|
| 526 |
+
%shift = shufflevector <2 x float> %460, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 527 |
+
%foldExtExtBinop = fadd <2 x float> %460, %shift, !dbg !60
|
| 528 |
+
%510 = extractelement <2 x float> %foldExtExtBinop, i64 0, !dbg !60
|
| 529 |
+
%shift66 = shufflevector <2 x float> %467, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 530 |
+
%foldExtExtBinop67 = fadd <2 x float> %467, %shift66, !dbg !60
|
| 531 |
+
%511 = extractelement <2 x float> %foldExtExtBinop67, i64 0, !dbg !60
|
| 532 |
+
%shift69 = shufflevector <2 x float> %474, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 533 |
+
%foldExtExtBinop70 = fadd <2 x float> %474, %shift69, !dbg !60
|
| 534 |
+
%512 = extractelement <2 x float> %foldExtExtBinop70, i64 0, !dbg !60
|
| 535 |
+
%shift72 = shufflevector <2 x float> %481, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 536 |
+
%foldExtExtBinop73 = fadd <2 x float> %481, %shift72, !dbg !60
|
| 537 |
+
%513 = extractelement <2 x float> %foldExtExtBinop73, i64 0, !dbg !60
|
| 538 |
+
%shift75 = shufflevector <2 x float> %488, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 539 |
+
%foldExtExtBinop76 = fadd <2 x float> %488, %shift75, !dbg !60
|
| 540 |
+
%514 = extractelement <2 x float> %foldExtExtBinop76, i64 0, !dbg !60
|
| 541 |
+
%shift78 = shufflevector <2 x float> %495, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 542 |
+
%foldExtExtBinop79 = fadd <2 x float> %495, %shift78, !dbg !60
|
| 543 |
+
%515 = extractelement <2 x float> %foldExtExtBinop79, i64 0, !dbg !60
|
| 544 |
+
%shift81 = shufflevector <2 x float> %502, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 545 |
+
%foldExtExtBinop82 = fadd <2 x float> %502, %shift81, !dbg !60
|
| 546 |
+
%516 = extractelement <2 x float> %foldExtExtBinop82, i64 0, !dbg !60
|
| 547 |
+
%shift84 = shufflevector <2 x float> %509, <2 x float> poison, <2 x i32> <i32 1, i32 poison>, !dbg !60
|
| 548 |
+
%foldExtExtBinop85 = fadd <2 x float> %509, %shift84, !dbg !60
|
| 549 |
+
%517 = extractelement <2 x float> %foldExtExtBinop85, i64 0, !dbg !60
|
| 550 |
+
%518 = getelementptr bfloat, ptr addrspace(1) %4, i64 %36, !dbg !61
|
| 551 |
+
%519 = getelementptr bfloat, ptr addrspace(1) %4, i64 %38, !dbg !61
|
| 552 |
+
%520 = getelementptr bfloat, ptr addrspace(1) %4, i64 %40, !dbg !61
|
| 553 |
+
%521 = getelementptr bfloat, ptr addrspace(1) %4, i64 %42, !dbg !61
|
| 554 |
+
%522 = getelementptr bfloat, ptr addrspace(1) %4, i64 %44, !dbg !61
|
| 555 |
+
%523 = getelementptr bfloat, ptr addrspace(1) %4, i64 %46, !dbg !61
|
| 556 |
+
%524 = getelementptr bfloat, ptr addrspace(1) %4, i64 %48, !dbg !61
|
| 557 |
+
%525 = getelementptr bfloat, ptr addrspace(1) %4, i64 %50, !dbg !61
|
| 558 |
+
%526 = fptrunc float %510 to bfloat, !dbg !62
|
| 559 |
+
%527 = fptrunc float %511 to bfloat, !dbg !62
|
| 560 |
+
%528 = fptrunc float %512 to bfloat, !dbg !62
|
| 561 |
+
%529 = fptrunc float %513 to bfloat, !dbg !62
|
| 562 |
+
%530 = fptrunc float %514 to bfloat, !dbg !62
|
| 563 |
+
%531 = fptrunc float %515 to bfloat, !dbg !62
|
| 564 |
+
%532 = fptrunc float %516 to bfloat, !dbg !62
|
| 565 |
+
%533 = fptrunc float %517 to bfloat, !dbg !62
|
| 566 |
+
%534 = bitcast bfloat %526 to i16, !dbg !62
|
| 567 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %534, ptr addrspace(1) %518, i1 %88) #4, !dbg !62
|
| 568 |
+
%535 = bitcast bfloat %527 to i16, !dbg !62
|
| 569 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %535, ptr addrspace(1) %519, i1 %91) #4, !dbg !62
|
| 570 |
+
%536 = bitcast bfloat %528 to i16, !dbg !62
|
| 571 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %536, ptr addrspace(1) %520, i1 %94) #4, !dbg !62
|
| 572 |
+
%537 = bitcast bfloat %529 to i16, !dbg !62
|
| 573 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %537, ptr addrspace(1) %521, i1 %97) #4, !dbg !62
|
| 574 |
+
%538 = bitcast bfloat %530 to i16, !dbg !62
|
| 575 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %538, ptr addrspace(1) %522, i1 %100) #4, !dbg !62
|
| 576 |
+
%539 = bitcast bfloat %531 to i16, !dbg !62
|
| 577 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %539, ptr addrspace(1) %523, i1 %103) #4, !dbg !62
|
| 578 |
+
%540 = bitcast bfloat %532 to i16, !dbg !62
|
| 579 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %540, ptr addrspace(1) %524, i1 %106) #4, !dbg !62
|
| 580 |
+
%541 = bitcast bfloat %533 to i16, !dbg !62
|
| 581 |
+
tail call void asm sideeffect "@$2 st.global.b16 [ $1 + 0 ], { $0 };", "c,l,b"(i16 %541, ptr addrspace(1) %525, i1 %109) #4, !dbg !62
|
| 582 |
+
ret void, !dbg !63
|
| 583 |
+
}
|
| 584 |
+
|
| 585 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 586 |
+
declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #2
|
| 587 |
+
|
| 588 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 589 |
+
declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #2
|
| 590 |
+
|
| 591 |
+
; Function Attrs: convergent nocallback nounwind
|
| 592 |
+
declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3
|
| 593 |
+
|
| 594 |
+
attributes #0 = { noreturn }
|
| 595 |
+
attributes #1 = { "nvvm.reqntid"="128" }
|
| 596 |
+
attributes #2 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
|
| 597 |
+
attributes #3 = { convergent nocallback nounwind }
|
| 598 |
+
attributes #4 = { nounwind }
|
| 599 |
+
|
| 600 |
+
!llvm.dbg.cu = !{!0}
|
| 601 |
+
!llvm.module.flags = !{!2, !3}
|
| 602 |
+
!llvm.ident = !{!4}
|
| 603 |
+
|
| 604 |
+
!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
|
| 605 |
+
!1 = !DIFile(filename: "crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv")
|
| 606 |
+
!2 = !{i32 2, !"Debug Info Version", i32 3}
|
| 607 |
+
!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
|
| 608 |
+
!4 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}
|
| 609 |
+
!5 = !DISubprogram(name: "__assertfail", linkageName: "__assertfail", scope: !6, file: !6, type: !7, spFlags: DISPFlagOptimized)
|
| 610 |
+
!6 = !DIFile(filename: "<unknown>", directory: "")
|
| 611 |
+
!7 = !DISubroutineType(cc: DW_CC_normal, types: !8)
|
| 612 |
+
!8 = !{}
|
| 613 |
+
!9 = distinct !DISubprogram(name: "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0", linkageName: "triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0", scope: !1, file: !1, line: 18, type: !7, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
|
| 614 |
+
!10 = !DILocation(line: 19, column: 28, scope: !9)
|
| 615 |
+
!11 = !DILocation(line: 19, column: 33, scope: !9)
|
| 616 |
+
!12 = !DILocation(line: 20, column: 36, scope: !9)
|
| 617 |
+
!13 = !DILocation(line: 20, column: 23, scope: !9)
|
| 618 |
+
!14 = !DILocation(line: 23, column: 21, scope: !9)
|
| 619 |
+
!15 = !DILocation(line: 26, column: 30, scope: !9)
|
| 620 |
+
!16 = !DILocation(line: 23, column: 28, scope: !9)
|
| 621 |
+
!17 = !DILocation(line: 26, column: 35, scope: !9)
|
| 622 |
+
!18 = !DILocation(line: 27, column: 30, scope: !9)
|
| 623 |
+
!19 = !DILocation(line: 21, column: 21, scope: !9)
|
| 624 |
+
!20 = !DILocation(line: 27, column: 35, scope: !9)
|
| 625 |
+
!21 = !DILocation(line: 30, column: 18, scope: !9)
|
| 626 |
+
!22 = !DILocation(line: 31, column: 32, scope: !9)
|
| 627 |
+
!23 = !DILocation(line: 32, column: 28, scope: !9)
|
| 628 |
+
!24 = !DILocation(line: 32, column: 44, scope: !9)
|
| 629 |
+
!25 = !DILocation(line: 32, column: 37, scope: !9)
|
| 630 |
+
!26 = !DILocation(line: 32, column: 52, scope: !9)
|
| 631 |
+
!27 = !DILocation(line: 32, column: 62, scope: !9)
|
| 632 |
+
!28 = !DILocation(line: 24, column: 19, scope: !9)
|
| 633 |
+
!29 = !DILocation(line: 33, column: 39, scope: !9)
|
| 634 |
+
!30 = !DILocation(line: 40, column: 35, scope: !9)
|
| 635 |
+
!31 = !DILocation(line: 33, column: 30, scope: !9)
|
| 636 |
+
!32 = !DILocation(line: 33, column: 46, scope: !9)
|
| 637 |
+
!33 = !DILocation(line: 38, column: 31, scope: !9)
|
| 638 |
+
!34 = !DILocation(line: 38, column: 18, scope: !9)
|
| 639 |
+
!35 = !DILocation(line: 39, column: 19, scope: !9)
|
| 640 |
+
!36 = !DILocation(line: 40, column: 31, scope: !9)
|
| 641 |
+
!37 = !DILocation(line: 40, column: 68, scope: !9)
|
| 642 |
+
!38 = !DILocation(line: 40, column: 60, scope: !9)
|
| 643 |
+
!39 = !DILocation(line: 44, column: 20, scope: !9)
|
| 644 |
+
!40 = !DILocation(line: 47, column: 47, scope: !9)
|
| 645 |
+
!41 = !DILocation(line: 47, column: 31, scope: !9)
|
| 646 |
+
!42 = !DILocation(line: 47, column: 81, scope: !9)
|
| 647 |
+
!43 = !DILocation(line: 47, column: 73, scope: !9)
|
| 648 |
+
!44 = !DILocation(line: 51, column: 34, scope: !9)
|
| 649 |
+
!45 = !DILocation(line: 52, column: 28, scope: !9)
|
| 650 |
+
!46 = !DILocation(line: 52, column: 46, scope: !9)
|
| 651 |
+
!47 = !DILocation(line: 52, column: 38, scope: !9)
|
| 652 |
+
!48 = !DILocation(line: 52, column: 54, scope: !9)
|
| 653 |
+
!49 = !DILocation(line: 52, column: 64, scope: !9)
|
| 654 |
+
!50 = !DILocation(line: 40, column: 119, scope: !9)
|
| 655 |
+
!51 = !DILocation(line: 41, column: 13, scope: !9)
|
| 656 |
+
!52 = !DILocation(line: 47, column: 132, scope: !9)
|
| 657 |
+
!53 = !DILocation(line: 0, scope: !9)
|
| 658 |
+
!54 = !DILocation(line: 26, column: 75, scope: !9)
|
| 659 |
+
!55 = !DILocation(line: 53, column: 40, scope: !9)
|
| 660 |
+
!56 = !DILocation(line: 53, column: 31, scope: !9)
|
| 661 |
+
!57 = !DILocation(line: 53, column: 48, scope: !9)
|
| 662 |
+
!58 = !DILocation(line: 33, column: 86, scope: !9)
|
| 663 |
+
!59 = !DILocation(line: 34, column: 18, scope: !9)
|
| 664 |
+
!60 = !DILocation(line: 55, column: 19, scope: !9)
|
| 665 |
+
!61 = !DILocation(line: 56, column: 25, scope: !9)
|
| 666 |
+
!62 = !DILocation(line: 56, column: 37, scope: !9)
|
| 667 |
+
!63 = !DILocation(line: 56, column: 4, scope: !9)
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ptx
ADDED
|
@@ -0,0 +1,1534 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
//
|
| 2 |
+
// Generated by LLVM NVPTX Back-End
|
| 3 |
+
//
|
| 4 |
+
|
| 5 |
+
.version 8.7
|
| 6 |
+
.target sm_90a
|
| 7 |
+
.address_size 64
|
| 8 |
+
|
| 9 |
+
// .globl triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0 // -- Begin function triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0
|
| 10 |
+
.extern .func __assertfail
|
| 11 |
+
(
|
| 12 |
+
.param .b64 __assertfail_param_0,
|
| 13 |
+
.param .b64 __assertfail_param_1,
|
| 14 |
+
.param .b32 __assertfail_param_2,
|
| 15 |
+
.param .b64 __assertfail_param_3,
|
| 16 |
+
.param .b64 __assertfail_param_4
|
| 17 |
+
)
|
| 18 |
+
.noreturn;
|
| 19 |
+
.global .align 1 .b8 assertFunc_1[8] = {117, 110, 107, 110, 111, 119, 110};
|
| 20 |
+
.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, 114, 118, 47, 99, 114, 118, 120, 98, 122, 116, 107, 114, 51, 55, 50, 102, 103, 100, 110, 55, 98, 103, 114, 117, 100, 50, 50, 115, 51, 119, 109, 100, 50, 105, 115, 105, 100, 119, 111, 52, 101, 107, 52, 104, 108, 100, 110, 53, 54, 116, 117, 118, 50, 100, 106, 46, 112, 121};
|
| 21 |
+
.global .align 1 .b8 assertMessage_1[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, 50, 53, 32, 60, 32, 107, 115, 52};
|
| 22 |
+
.global .align 1 .b8 assertFunc_0[8] = {117, 110, 107, 110, 111, 119, 110};
|
| 23 |
+
.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, 114, 118, 47, 99, 114, 118, 120, 98, 122, 116, 107, 114, 51, 55, 50, 102, 103, 100, 110, 55, 98, 103, 114, 117, 100, 50, 50, 115, 51, 119, 109, 100, 50, 105, 115, 105, 100, 119, 111, 52, 101, 107, 52, 104, 108, 100, 110, 53, 54, 116, 117, 118, 50, 100, 106, 46, 112, 121};
|
| 24 |
+
.global .align 1 .b8 assertMessage_0[37] = {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, 53, 32, 60, 32, 107, 115, 50};
|
| 25 |
+
// @triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0
|
| 26 |
+
.visible .entry triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0(
|
| 27 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_0,
|
| 28 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_1,
|
| 29 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_2,
|
| 30 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_3,
|
| 31 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_4,
|
| 32 |
+
.param .u64 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_5,
|
| 33 |
+
.param .u64 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_6,
|
| 34 |
+
.param .u64 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_7,
|
| 35 |
+
.param .u64 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_8,
|
| 36 |
+
.param .u64 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_9,
|
| 37 |
+
.param .u32 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_10,
|
| 38 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_11,
|
| 39 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_12
|
| 40 |
+
)
|
| 41 |
+
.reqntid 128
|
| 42 |
+
{
|
| 43 |
+
.reg .pred %p<179>;
|
| 44 |
+
.reg .b16 %rs<165>;
|
| 45 |
+
.reg .b32 %r<160>;
|
| 46 |
+
.reg .b64 %rd<500>;
|
| 47 |
+
.loc 1 18 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:18:0
|
| 48 |
+
$L__func_begin0:
|
| 49 |
+
.loc 1 18 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:18:0
|
| 50 |
+
|
| 51 |
+
// %bb.0:
|
| 52 |
+
ld.param.b64 %rd103, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_5];
|
| 53 |
+
$L__tmp0:
|
| 54 |
+
.loc 1 19 28 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:19:28
|
| 55 |
+
mov.u32 %r26, %ctaid.x;
|
| 56 |
+
.loc 1 19 33 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:19:33
|
| 57 |
+
shl.b32 %r27, %r26, 10;
|
| 58 |
+
.loc 1 20 36 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:20:36
|
| 59 |
+
mov.u32 %r28, %tid.x;
|
| 60 |
+
shl.b32 %r29, %r28, 3;
|
| 61 |
+
and.b32 %r30, %r29, 1016;
|
| 62 |
+
.loc 1 20 23 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:20:23
|
| 63 |
+
or.b32 %r1, %r30, %r27;
|
| 64 |
+
or.b32 %r2, %r1, 1;
|
| 65 |
+
or.b32 %r3, %r1, 2;
|
| 66 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 67 |
+
cvt.s64.s32 %rd7, %r2;
|
| 68 |
+
cvt.s64.s32 %rd8, %r1;
|
| 69 |
+
or.b64 %rd108, %rd8, %rd103;
|
| 70 |
+
and.b64 %rd109, %rd108, -4294967296;
|
| 71 |
+
setp.ne.b64 %p9, %rd109, 0;
|
| 72 |
+
@%p9 bra $L__BB0_2;
|
| 73 |
+
bra.uni $L__BB0_1;
|
| 74 |
+
$L__BB0_2:
|
| 75 |
+
div.s64 %rd484, %rd8, %rd103;
|
| 76 |
+
bra.uni $L__BB0_3;
|
| 77 |
+
$L__BB0_1:
|
| 78 |
+
cvt.u32.u64 %r31, %rd103;
|
| 79 |
+
cvt.u32.u64 %r32, %rd8;
|
| 80 |
+
div.u32 %r33, %r32, %r31;
|
| 81 |
+
cvt.u64.u32 %rd484, %r33;
|
| 82 |
+
$L__BB0_3:
|
| 83 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 84 |
+
or.b32 %r4, %r1, 3;
|
| 85 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 86 |
+
cvt.s64.s32 %rd6, %r3;
|
| 87 |
+
or.b64 %rd110, %rd7, %rd103;
|
| 88 |
+
and.b64 %rd111, %rd110, -4294967296;
|
| 89 |
+
setp.ne.b64 %p10, %rd111, 0;
|
| 90 |
+
@%p10 bra $L__BB0_5;
|
| 91 |
+
bra.uni $L__BB0_4;
|
| 92 |
+
$L__BB0_5:
|
| 93 |
+
div.s64 %rd485, %rd7, %rd103;
|
| 94 |
+
bra.uni $L__BB0_6;
|
| 95 |
+
$L__BB0_4:
|
| 96 |
+
cvt.u32.u64 %r34, %rd103;
|
| 97 |
+
cvt.u32.u64 %r35, %rd7;
|
| 98 |
+
div.u32 %r36, %r35, %r34;
|
| 99 |
+
cvt.u64.u32 %rd485, %r36;
|
| 100 |
+
$L__BB0_6:
|
| 101 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 102 |
+
or.b32 %r5, %r1, 4;
|
| 103 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 104 |
+
cvt.s64.s32 %rd5, %r4;
|
| 105 |
+
or.b64 %rd112, %rd6, %rd103;
|
| 106 |
+
and.b64 %rd113, %rd112, -4294967296;
|
| 107 |
+
setp.ne.b64 %p11, %rd113, 0;
|
| 108 |
+
@%p11 bra $L__BB0_8;
|
| 109 |
+
bra.uni $L__BB0_7;
|
| 110 |
+
$L__BB0_8:
|
| 111 |
+
div.s64 %rd486, %rd6, %rd103;
|
| 112 |
+
bra.uni $L__BB0_9;
|
| 113 |
+
$L__BB0_7:
|
| 114 |
+
cvt.u32.u64 %r37, %rd103;
|
| 115 |
+
cvt.u32.u64 %r38, %rd6;
|
| 116 |
+
div.u32 %r39, %r38, %r37;
|
| 117 |
+
cvt.u64.u32 %rd486, %r39;
|
| 118 |
+
$L__BB0_9:
|
| 119 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 120 |
+
or.b32 %r6, %r1, 5;
|
| 121 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 122 |
+
cvt.s64.s32 %rd4, %r5;
|
| 123 |
+
or.b64 %rd114, %rd5, %rd103;
|
| 124 |
+
and.b64 %rd115, %rd114, -4294967296;
|
| 125 |
+
setp.ne.b64 %p12, %rd115, 0;
|
| 126 |
+
@%p12 bra $L__BB0_11;
|
| 127 |
+
bra.uni $L__BB0_10;
|
| 128 |
+
$L__BB0_11:
|
| 129 |
+
div.s64 %rd487, %rd5, %rd103;
|
| 130 |
+
bra.uni $L__BB0_12;
|
| 131 |
+
$L__BB0_10:
|
| 132 |
+
cvt.u32.u64 %r40, %rd103;
|
| 133 |
+
cvt.u32.u64 %r41, %rd5;
|
| 134 |
+
div.u32 %r42, %r41, %r40;
|
| 135 |
+
cvt.u64.u32 %rd487, %r42;
|
| 136 |
+
$L__BB0_12:
|
| 137 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 138 |
+
or.b32 %r7, %r1, 6;
|
| 139 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 140 |
+
cvt.s64.s32 %rd3, %r6;
|
| 141 |
+
or.b64 %rd116, %rd4, %rd103;
|
| 142 |
+
and.b64 %rd117, %rd116, -4294967296;
|
| 143 |
+
setp.ne.b64 %p13, %rd117, 0;
|
| 144 |
+
@%p13 bra $L__BB0_14;
|
| 145 |
+
bra.uni $L__BB0_13;
|
| 146 |
+
$L__BB0_14:
|
| 147 |
+
div.s64 %rd488, %rd4, %rd103;
|
| 148 |
+
bra.uni $L__BB0_15;
|
| 149 |
+
$L__BB0_13:
|
| 150 |
+
cvt.u32.u64 %r43, %rd103;
|
| 151 |
+
cvt.u32.u64 %r44, %rd4;
|
| 152 |
+
div.u32 %r45, %r44, %r43;
|
| 153 |
+
cvt.u64.u32 %rd488, %r45;
|
| 154 |
+
$L__BB0_15:
|
| 155 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 156 |
+
or.b32 %r8, %r1, 7;
|
| 157 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 158 |
+
cvt.s64.s32 %rd2, %r7;
|
| 159 |
+
or.b64 %rd118, %rd3, %rd103;
|
| 160 |
+
and.b64 %rd119, %rd118, -4294967296;
|
| 161 |
+
setp.ne.b64 %p14, %rd119, 0;
|
| 162 |
+
@%p14 bra $L__BB0_17;
|
| 163 |
+
bra.uni $L__BB0_16;
|
| 164 |
+
$L__BB0_17:
|
| 165 |
+
div.s64 %rd489, %rd3, %rd103;
|
| 166 |
+
bra.uni $L__BB0_18;
|
| 167 |
+
$L__BB0_16:
|
| 168 |
+
cvt.u32.u64 %r46, %rd103;
|
| 169 |
+
cvt.u32.u64 %r47, %rd3;
|
| 170 |
+
div.u32 %r48, %r47, %r46;
|
| 171 |
+
cvt.u64.u32 %rd489, %r48;
|
| 172 |
+
$L__BB0_18:
|
| 173 |
+
cvt.s64.s32 %rd1, %r8;
|
| 174 |
+
or.b64 %rd120, %rd2, %rd103;
|
| 175 |
+
and.b64 %rd121, %rd120, -4294967296;
|
| 176 |
+
setp.ne.b64 %p15, %rd121, 0;
|
| 177 |
+
@%p15 bra $L__BB0_20;
|
| 178 |
+
bra.uni $L__BB0_19;
|
| 179 |
+
$L__BB0_20:
|
| 180 |
+
div.s64 %rd490, %rd2, %rd103;
|
| 181 |
+
bra.uni $L__BB0_21;
|
| 182 |
+
$L__BB0_19:
|
| 183 |
+
cvt.u32.u64 %r49, %rd103;
|
| 184 |
+
cvt.u32.u64 %r50, %rd2;
|
| 185 |
+
div.u32 %r51, %r50, %r49;
|
| 186 |
+
cvt.u64.u32 %rd490, %r51;
|
| 187 |
+
$L__BB0_21:
|
| 188 |
+
.loc 1 0 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0:21
|
| 189 |
+
ld.param.b64 %rd104, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_6];
|
| 190 |
+
.loc 1 23 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:21
|
| 191 |
+
or.b64 %rd122, %rd1, %rd103;
|
| 192 |
+
and.b64 %rd123, %rd122, -4294967296;
|
| 193 |
+
setp.ne.b64 %p16, %rd123, 0;
|
| 194 |
+
@%p16 bra $L__BB0_23;
|
| 195 |
+
bra.uni $L__BB0_22;
|
| 196 |
+
$L__BB0_23:
|
| 197 |
+
div.s64 %rd491, %rd1, %rd103;
|
| 198 |
+
bra.uni $L__BB0_24;
|
| 199 |
+
$L__BB0_22:
|
| 200 |
+
cvt.u32.u64 %r52, %rd103;
|
| 201 |
+
cvt.u32.u64 %r53, %rd1;
|
| 202 |
+
div.u32 %r54, %r53, %r52;
|
| 203 |
+
cvt.u64.u32 %rd491, %r54;
|
| 204 |
+
$L__BB0_24:
|
| 205 |
+
.loc 1 23 28 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:28
|
| 206 |
+
or.b64 %rd124, %rd484, %rd104;
|
| 207 |
+
and.b64 %rd125, %rd124, -4294967296;
|
| 208 |
+
setp.ne.b64 %p17, %rd125, 0;
|
| 209 |
+
@%p17 bra $L__BB0_26;
|
| 210 |
+
bra.uni $L__BB0_25;
|
| 211 |
+
$L__BB0_26:
|
| 212 |
+
rem.s64 %rd492, %rd484, %rd104;
|
| 213 |
+
bra.uni $L__BB0_27;
|
| 214 |
+
$L__BB0_25:
|
| 215 |
+
cvt.u32.u64 %r55, %rd104;
|
| 216 |
+
cvt.u32.u64 %r56, %rd484;
|
| 217 |
+
rem.u32 %r57, %r56, %r55;
|
| 218 |
+
cvt.u64.u32 %rd492, %r57;
|
| 219 |
+
$L__BB0_27:
|
| 220 |
+
or.b64 %rd126, %rd485, %rd104;
|
| 221 |
+
and.b64 %rd127, %rd126, -4294967296;
|
| 222 |
+
setp.ne.b64 %p18, %rd127, 0;
|
| 223 |
+
@%p18 bra $L__BB0_29;
|
| 224 |
+
bra.uni $L__BB0_28;
|
| 225 |
+
$L__BB0_29:
|
| 226 |
+
rem.s64 %rd493, %rd485, %rd104;
|
| 227 |
+
bra.uni $L__BB0_30;
|
| 228 |
+
$L__BB0_28:
|
| 229 |
+
cvt.u32.u64 %r58, %rd104;
|
| 230 |
+
cvt.u32.u64 %r59, %rd485;
|
| 231 |
+
rem.u32 %r60, %r59, %r58;
|
| 232 |
+
cvt.u64.u32 %rd493, %r60;
|
| 233 |
+
$L__BB0_30:
|
| 234 |
+
or.b64 %rd128, %rd486, %rd104;
|
| 235 |
+
and.b64 %rd129, %rd128, -4294967296;
|
| 236 |
+
setp.ne.b64 %p19, %rd129, 0;
|
| 237 |
+
@%p19 bra $L__BB0_32;
|
| 238 |
+
bra.uni $L__BB0_31;
|
| 239 |
+
$L__BB0_32:
|
| 240 |
+
rem.s64 %rd494, %rd486, %rd104;
|
| 241 |
+
bra.uni $L__BB0_33;
|
| 242 |
+
$L__BB0_31:
|
| 243 |
+
cvt.u32.u64 %r61, %rd104;
|
| 244 |
+
cvt.u32.u64 %r62, %rd486;
|
| 245 |
+
rem.u32 %r63, %r62, %r61;
|
| 246 |
+
cvt.u64.u32 %rd494, %r63;
|
| 247 |
+
$L__BB0_33:
|
| 248 |
+
or.b64 %rd130, %rd487, %rd104;
|
| 249 |
+
and.b64 %rd131, %rd130, -4294967296;
|
| 250 |
+
setp.ne.b64 %p20, %rd131, 0;
|
| 251 |
+
@%p20 bra $L__BB0_35;
|
| 252 |
+
bra.uni $L__BB0_34;
|
| 253 |
+
$L__BB0_35:
|
| 254 |
+
rem.s64 %rd495, %rd487, %rd104;
|
| 255 |
+
bra.uni $L__BB0_36;
|
| 256 |
+
$L__BB0_34:
|
| 257 |
+
cvt.u32.u64 %r64, %rd104;
|
| 258 |
+
cvt.u32.u64 %r65, %rd487;
|
| 259 |
+
rem.u32 %r66, %r65, %r64;
|
| 260 |
+
cvt.u64.u32 %rd495, %r66;
|
| 261 |
+
$L__BB0_36:
|
| 262 |
+
or.b64 %rd132, %rd488, %rd104;
|
| 263 |
+
and.b64 %rd133, %rd132, -4294967296;
|
| 264 |
+
setp.ne.b64 %p21, %rd133, 0;
|
| 265 |
+
@%p21 bra $L__BB0_38;
|
| 266 |
+
bra.uni $L__BB0_37;
|
| 267 |
+
$L__BB0_38:
|
| 268 |
+
rem.s64 %rd496, %rd488, %rd104;
|
| 269 |
+
bra.uni $L__BB0_39;
|
| 270 |
+
$L__BB0_37:
|
| 271 |
+
cvt.u32.u64 %r67, %rd104;
|
| 272 |
+
cvt.u32.u64 %r68, %rd488;
|
| 273 |
+
rem.u32 %r69, %r68, %r67;
|
| 274 |
+
cvt.u64.u32 %rd496, %r69;
|
| 275 |
+
$L__BB0_39:
|
| 276 |
+
or.b64 %rd134, %rd489, %rd104;
|
| 277 |
+
and.b64 %rd135, %rd134, -4294967296;
|
| 278 |
+
setp.ne.b64 %p22, %rd135, 0;
|
| 279 |
+
@%p22 bra $L__BB0_41;
|
| 280 |
+
bra.uni $L__BB0_40;
|
| 281 |
+
$L__BB0_41:
|
| 282 |
+
rem.s64 %rd497, %rd489, %rd104;
|
| 283 |
+
bra.uni $L__BB0_42;
|
| 284 |
+
$L__BB0_40:
|
| 285 |
+
cvt.u32.u64 %r70, %rd104;
|
| 286 |
+
cvt.u32.u64 %r71, %rd489;
|
| 287 |
+
rem.u32 %r72, %r71, %r70;
|
| 288 |
+
cvt.u64.u32 %rd497, %r72;
|
| 289 |
+
$L__BB0_42:
|
| 290 |
+
or.b64 %rd136, %rd490, %rd104;
|
| 291 |
+
and.b64 %rd137, %rd136, -4294967296;
|
| 292 |
+
setp.ne.b64 %p23, %rd137, 0;
|
| 293 |
+
@%p23 bra $L__BB0_44;
|
| 294 |
+
bra.uni $L__BB0_43;
|
| 295 |
+
$L__BB0_44:
|
| 296 |
+
rem.s64 %rd498, %rd490, %rd104;
|
| 297 |
+
bra.uni $L__BB0_45;
|
| 298 |
+
$L__BB0_43:
|
| 299 |
+
cvt.u32.u64 %r73, %rd104;
|
| 300 |
+
cvt.u32.u64 %r74, %rd490;
|
| 301 |
+
rem.u32 %r75, %r74, %r73;
|
| 302 |
+
cvt.u64.u32 %rd498, %r75;
|
| 303 |
+
$L__BB0_45:
|
| 304 |
+
.loc 1 0 28 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0:28
|
| 305 |
+
ld.param.b32 %r25, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_10];
|
| 306 |
+
ld.param.b64 %rd105, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_7];
|
| 307 |
+
ld.param.b64 %rd99, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_1];
|
| 308 |
+
ld.param.b64 %rd98, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_0];
|
| 309 |
+
.loc 1 23 28 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:23:28
|
| 310 |
+
or.b64 %rd138, %rd491, %rd104;
|
| 311 |
+
and.b64 %rd139, %rd138, -4294967296;
|
| 312 |
+
setp.ne.b64 %p24, %rd139, 0;
|
| 313 |
+
@%p24 bra $L__BB0_47;
|
| 314 |
+
bra.uni $L__BB0_46;
|
| 315 |
+
$L__BB0_47:
|
| 316 |
+
rem.s64 %rd499, %rd491, %rd104;
|
| 317 |
+
bra.uni $L__BB0_48;
|
| 318 |
+
$L__BB0_46:
|
| 319 |
+
cvt.u32.u64 %r76, %rd104;
|
| 320 |
+
cvt.u32.u64 %r77, %rd491;
|
| 321 |
+
rem.u32 %r78, %r77, %r76;
|
| 322 |
+
cvt.u64.u32 %rd499, %r78;
|
| 323 |
+
$L__BB0_48:
|
| 324 |
+
.loc 1 26 30 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:26:30
|
| 325 |
+
shl.b64 %rd196, %rd8, 1;
|
| 326 |
+
add.s64 %rd141, %rd98, %rd196;
|
| 327 |
+
shl.b64 %rd197, %rd7, 1;
|
| 328 |
+
add.s64 %rd144, %rd98, %rd197;
|
| 329 |
+
shl.b64 %rd198, %rd6, 1;
|
| 330 |
+
add.s64 %rd147, %rd98, %rd198;
|
| 331 |
+
shl.b64 %rd199, %rd5, 1;
|
| 332 |
+
add.s64 %rd150, %rd98, %rd199;
|
| 333 |
+
shl.b64 %rd200, %rd4, 1;
|
| 334 |
+
add.s64 %rd153, %rd98, %rd200;
|
| 335 |
+
shl.b64 %rd201, %rd3, 1;
|
| 336 |
+
add.s64 %rd156, %rd98, %rd201;
|
| 337 |
+
shl.b64 %rd202, %rd2, 1;
|
| 338 |
+
add.s64 %rd159, %rd98, %rd202;
|
| 339 |
+
shl.b64 %rd203, %rd1, 1;
|
| 340 |
+
add.s64 %rd162, %rd98, %rd203;
|
| 341 |
+
.loc 1 26 35 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:26:35
|
| 342 |
+
// begin inline asm
|
| 343 |
+
mov.u64 %rd140, 0x0;
|
| 344 |
+
createpolicy.fractional.L2::evict_last.b64 %rd140, 1.0;
|
| 345 |
+
// end inline asm
|
| 346 |
+
.loc 1 27 30 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:27:30
|
| 347 |
+
shl.b64 %rd204, %rd492, 3;
|
| 348 |
+
add.s64 %rd166, %rd99, %rd204;
|
| 349 |
+
shl.b64 %rd205, %rd493, 3;
|
| 350 |
+
add.s64 %rd170, %rd99, %rd205;
|
| 351 |
+
shl.b64 %rd206, %rd494, 3;
|
| 352 |
+
add.s64 %rd174, %rd99, %rd206;
|
| 353 |
+
shl.b64 %rd207, %rd495, 3;
|
| 354 |
+
add.s64 %rd178, %rd99, %rd207;
|
| 355 |
+
shl.b64 %rd208, %rd496, 3;
|
| 356 |
+
add.s64 %rd182, %rd99, %rd208;
|
| 357 |
+
shl.b64 %rd209, %rd497, 3;
|
| 358 |
+
add.s64 %rd186, %rd99, %rd209;
|
| 359 |
+
shl.b64 %rd210, %rd498, 3;
|
| 360 |
+
add.s64 %rd190, %rd99, %rd210;
|
| 361 |
+
shl.b64 %rd211, %rd499, 3;
|
| 362 |
+
add.s64 %rd194, %rd99, %rd211;
|
| 363 |
+
.loc 1 21 21 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:21:21
|
| 364 |
+
setp.lt.s32 %p8, %r8, %r25;
|
| 365 |
+
setp.lt.s32 %p7, %r7, %r25;
|
| 366 |
+
setp.lt.s32 %p6, %r6, %r25;
|
| 367 |
+
setp.lt.s32 %p5, %r5, %r25;
|
| 368 |
+
setp.lt.s32 %p4, %r4, %r25;
|
| 369 |
+
setp.lt.s32 %p3, %r3, %r25;
|
| 370 |
+
setp.lt.s32 %p2, %r2, %r25;
|
| 371 |
+
setp.lt.s32 %p1, %r1, %r25;
|
| 372 |
+
.loc 1 26 35 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:26:35
|
| 373 |
+
// begin inline asm
|
| 374 |
+
mov.u16 %rs33, 0x0;
|
| 375 |
+
@%p1 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs33 }, [ %rd141 + 0 ], %rd140;
|
| 376 |
+
// end inline asm
|
| 377 |
+
// begin inline asm
|
| 378 |
+
mov.u64 %rd143, 0x0;
|
| 379 |
+
createpolicy.fractional.L2::evict_last.b64 %rd143, 1.0;
|
| 380 |
+
// end inline asm
|
| 381 |
+
// begin inline asm
|
| 382 |
+
mov.u16 %rs34, 0x0;
|
| 383 |
+
@%p2 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs34 }, [ %rd144 + 0 ], %rd143;
|
| 384 |
+
// end inline asm
|
| 385 |
+
// begin inline asm
|
| 386 |
+
mov.u64 %rd146, 0x0;
|
| 387 |
+
createpolicy.fractional.L2::evict_last.b64 %rd146, 1.0;
|
| 388 |
+
// end inline asm
|
| 389 |
+
// begin inline asm
|
| 390 |
+
mov.u16 %rs35, 0x0;
|
| 391 |
+
@%p3 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs35 }, [ %rd147 + 0 ], %rd146;
|
| 392 |
+
// end inline asm
|
| 393 |
+
// begin inline asm
|
| 394 |
+
mov.u64 %rd149, 0x0;
|
| 395 |
+
createpolicy.fractional.L2::evict_last.b64 %rd149, 1.0;
|
| 396 |
+
// end inline asm
|
| 397 |
+
// begin inline asm
|
| 398 |
+
mov.u16 %rs36, 0x0;
|
| 399 |
+
@%p4 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs36 }, [ %rd150 + 0 ], %rd149;
|
| 400 |
+
// end inline asm
|
| 401 |
+
// begin inline asm
|
| 402 |
+
mov.u64 %rd152, 0x0;
|
| 403 |
+
createpolicy.fractional.L2::evict_last.b64 %rd152, 1.0;
|
| 404 |
+
// end inline asm
|
| 405 |
+
// begin inline asm
|
| 406 |
+
mov.u16 %rs37, 0x0;
|
| 407 |
+
@%p5 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs37 }, [ %rd153 + 0 ], %rd152;
|
| 408 |
+
// end inline asm
|
| 409 |
+
// begin inline asm
|
| 410 |
+
mov.u64 %rd155, 0x0;
|
| 411 |
+
createpolicy.fractional.L2::evict_last.b64 %rd155, 1.0;
|
| 412 |
+
// end inline asm
|
| 413 |
+
// begin inline asm
|
| 414 |
+
mov.u16 %rs38, 0x0;
|
| 415 |
+
@%p6 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs38 }, [ %rd156 + 0 ], %rd155;
|
| 416 |
+
// end inline asm
|
| 417 |
+
// begin inline asm
|
| 418 |
+
mov.u64 %rd158, 0x0;
|
| 419 |
+
createpolicy.fractional.L2::evict_last.b64 %rd158, 1.0;
|
| 420 |
+
// end inline asm
|
| 421 |
+
// begin inline asm
|
| 422 |
+
mov.u16 %rs39, 0x0;
|
| 423 |
+
@%p7 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs39 }, [ %rd159 + 0 ], %rd158;
|
| 424 |
+
// end inline asm
|
| 425 |
+
// begin inline asm
|
| 426 |
+
mov.u64 %rd161, 0x0;
|
| 427 |
+
createpolicy.fractional.L2::evict_last.b64 %rd161, 1.0;
|
| 428 |
+
// end inline asm
|
| 429 |
+
// begin inline asm
|
| 430 |
+
mov.u16 %rs40, 0x0;
|
| 431 |
+
@%p8 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs40 }, [ %rd162 + 0 ], %rd161;
|
| 432 |
+
// end inline asm
|
| 433 |
+
.loc 1 27 35 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:27:35
|
| 434 |
+
// begin inline asm
|
| 435 |
+
mov.u64 %rd164, 0x0;
|
| 436 |
+
createpolicy.fractional.L2::evict_last.b64 %rd164, 1.0;
|
| 437 |
+
// end inline asm
|
| 438 |
+
// begin inline asm
|
| 439 |
+
mov.u64 %rd165, 0x0;
|
| 440 |
+
@%p1 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd165 }, [ %rd166 + 0 ], %rd164;
|
| 441 |
+
// end inline asm
|
| 442 |
+
// begin inline asm
|
| 443 |
+
mov.u64 %rd168, 0x0;
|
| 444 |
+
createpolicy.fractional.L2::evict_last.b64 %rd168, 1.0;
|
| 445 |
+
// end inline asm
|
| 446 |
+
// begin inline asm
|
| 447 |
+
mov.u64 %rd169, 0x0;
|
| 448 |
+
@%p2 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd169 }, [ %rd170 + 0 ], %rd168;
|
| 449 |
+
// end inline asm
|
| 450 |
+
// begin inline asm
|
| 451 |
+
mov.u64 %rd172, 0x0;
|
| 452 |
+
createpolicy.fractional.L2::evict_last.b64 %rd172, 1.0;
|
| 453 |
+
// end inline asm
|
| 454 |
+
// begin inline asm
|
| 455 |
+
mov.u64 %rd173, 0x0;
|
| 456 |
+
@%p3 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd173 }, [ %rd174 + 0 ], %rd172;
|
| 457 |
+
// end inline asm
|
| 458 |
+
// begin inline asm
|
| 459 |
+
mov.u64 %rd176, 0x0;
|
| 460 |
+
createpolicy.fractional.L2::evict_last.b64 %rd176, 1.0;
|
| 461 |
+
// end inline asm
|
| 462 |
+
// begin inline asm
|
| 463 |
+
mov.u64 %rd177, 0x0;
|
| 464 |
+
@%p4 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd177 }, [ %rd178 + 0 ], %rd176;
|
| 465 |
+
// end inline asm
|
| 466 |
+
// begin inline asm
|
| 467 |
+
mov.u64 %rd180, 0x0;
|
| 468 |
+
createpolicy.fractional.L2::evict_last.b64 %rd180, 1.0;
|
| 469 |
+
// end inline asm
|
| 470 |
+
// begin inline asm
|
| 471 |
+
mov.u64 %rd181, 0x0;
|
| 472 |
+
@%p5 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd181 }, [ %rd182 + 0 ], %rd180;
|
| 473 |
+
// end inline asm
|
| 474 |
+
// begin inline asm
|
| 475 |
+
mov.u64 %rd184, 0x0;
|
| 476 |
+
createpolicy.fractional.L2::evict_last.b64 %rd184, 1.0;
|
| 477 |
+
// end inline asm
|
| 478 |
+
// begin inline asm
|
| 479 |
+
mov.u64 %rd185, 0x0;
|
| 480 |
+
@%p6 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd185 }, [ %rd186 + 0 ], %rd184;
|
| 481 |
+
// end inline asm
|
| 482 |
+
// begin inline asm
|
| 483 |
+
mov.u64 %rd188, 0x0;
|
| 484 |
+
createpolicy.fractional.L2::evict_last.b64 %rd188, 1.0;
|
| 485 |
+
// end inline asm
|
| 486 |
+
// begin inline asm
|
| 487 |
+
mov.u64 %rd189, 0x0;
|
| 488 |
+
@%p7 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd189 }, [ %rd190 + 0 ], %rd188;
|
| 489 |
+
// end inline asm
|
| 490 |
+
// begin inline asm
|
| 491 |
+
mov.u64 %rd192, 0x0;
|
| 492 |
+
createpolicy.fractional.L2::evict_last.b64 %rd192, 1.0;
|
| 493 |
+
// end inline asm
|
| 494 |
+
// begin inline asm
|
| 495 |
+
mov.u64 %rd193, 0x0;
|
| 496 |
+
@%p8 ld.global.L1::evict_last.L2::cache_hint.b64 { %rd193 }, [ %rd194 + 0 ], %rd192;
|
| 497 |
+
// end inline asm
|
| 498 |
+
.loc 1 31 32 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:31:32
|
| 499 |
+
shr.s64 %rd212, %rd173, 63;
|
| 500 |
+
and.b64 %rd213, %rd212, %rd105;
|
| 501 |
+
shr.s64 %rd214, %rd177, 63;
|
| 502 |
+
and.b64 %rd215, %rd214, %rd105;
|
| 503 |
+
shr.s64 %rd216, %rd165, 63;
|
| 504 |
+
and.b64 %rd217, %rd216, %rd105;
|
| 505 |
+
shr.s64 %rd218, %rd169, 63;
|
| 506 |
+
and.b64 %rd219, %rd218, %rd105;
|
| 507 |
+
shr.s64 %rd220, %rd189, 63;
|
| 508 |
+
and.b64 %rd221, %rd220, %rd105;
|
| 509 |
+
shr.s64 %rd222, %rd193, 63;
|
| 510 |
+
and.b64 %rd223, %rd222, %rd105;
|
| 511 |
+
shr.s64 %rd224, %rd181, 63;
|
| 512 |
+
and.b64 %rd225, %rd224, %rd105;
|
| 513 |
+
shr.s64 %rd226, %rd185, 63;
|
| 514 |
+
and.b64 %rd227, %rd226, %rd105;
|
| 515 |
+
add.s64 %rd78, %rd227, %rd185;
|
| 516 |
+
add.s64 %rd77, %rd225, %rd181;
|
| 517 |
+
add.s64 %rd80, %rd223, %rd193;
|
| 518 |
+
add.s64 %rd79, %rd221, %rd189;
|
| 519 |
+
add.s64 %rd74, %rd219, %rd169;
|
| 520 |
+
add.s64 %rd73, %rd217, %rd165;
|
| 521 |
+
add.s64 %rd76, %rd215, %rd177;
|
| 522 |
+
add.s64 %rd75, %rd213, %rd173;
|
| 523 |
+
.loc 1 32 28 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:28
|
| 524 |
+
setp.lt.s64 %p41, %rd75, 0;
|
| 525 |
+
setp.lt.s64 %p42, %rd76, 0;
|
| 526 |
+
setp.lt.s64 %p43, %rd73, 0;
|
| 527 |
+
setp.lt.s64 %p44, %rd74, 0;
|
| 528 |
+
setp.lt.s64 %p45, %rd79, 0;
|
| 529 |
+
setp.lt.s64 %p46, %rd80, 0;
|
| 530 |
+
setp.lt.s64 %p47, %rd77, 0;
|
| 531 |
+
setp.lt.s64 %p48, %rd78, 0;
|
| 532 |
+
.loc 1 32 44 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:44
|
| 533 |
+
setp.ge.s64 %p49, %rd75, %rd105;
|
| 534 |
+
setp.ge.s64 %p50, %rd76, %rd105;
|
| 535 |
+
setp.ge.s64 %p51, %rd73, %rd105;
|
| 536 |
+
setp.ge.s64 %p52, %rd74, %rd105;
|
| 537 |
+
setp.ge.s64 %p53, %rd79, %rd105;
|
| 538 |
+
setp.ge.s64 %p54, %rd80, %rd105;
|
| 539 |
+
setp.ge.s64 %p55, %rd77, %rd105;
|
| 540 |
+
setp.ge.s64 %p56, %rd78, %rd105;
|
| 541 |
+
.loc 1 32 37 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:37
|
| 542 |
+
or.pred %p57, %p48, %p56;
|
| 543 |
+
or.pred %p58, %p47, %p55;
|
| 544 |
+
or.pred %p59, %p46, %p54;
|
| 545 |
+
or.pred %p60, %p45, %p53;
|
| 546 |
+
or.pred %p61, %p44, %p52;
|
| 547 |
+
or.pred %p62, %p43, %p51;
|
| 548 |
+
or.pred %p63, %p42, %p50;
|
| 549 |
+
or.pred %p64, %p41, %p49;
|
| 550 |
+
.loc 1 32 52 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:52
|
| 551 |
+
and.pred %p65, %p3, %p64;
|
| 552 |
+
selp.b16 %rs41, 1, 0, %p65;
|
| 553 |
+
shl.b16 %rs42, %rs41, 2;
|
| 554 |
+
and.pred %p66, %p4, %p63;
|
| 555 |
+
selp.b16 %rs43, -1, 0, %p66;
|
| 556 |
+
shl.b16 %rs44, %rs43, 3;
|
| 557 |
+
or.b16 %rs45, %rs44, %rs42;
|
| 558 |
+
and.pred %p67, %p1, %p62;
|
| 559 |
+
selp.b16 %rs46, 1, 0, %p67;
|
| 560 |
+
and.pred %p68, %p2, %p61;
|
| 561 |
+
selp.b16 %rs47, -1, 0, %p68;
|
| 562 |
+
shl.b16 %rs48, %rs47, 1;
|
| 563 |
+
or.b16 %rs49, %rs46, %rs48;
|
| 564 |
+
and.b16 %rs50, %rs49, 3;
|
| 565 |
+
or.b16 %rs51, %rs50, %rs45;
|
| 566 |
+
and.b16 %rs52, %rs51, 15;
|
| 567 |
+
and.pred %p69, %p7, %p60;
|
| 568 |
+
selp.b16 %rs53, 1, 0, %p69;
|
| 569 |
+
shl.b16 %rs54, %rs53, 2;
|
| 570 |
+
and.pred %p70, %p8, %p59;
|
| 571 |
+
selp.b16 %rs55, -1, 0, %p70;
|
| 572 |
+
shl.b16 %rs56, %rs55, 3;
|
| 573 |
+
or.b16 %rs57, %rs56, %rs54;
|
| 574 |
+
and.pred %p71, %p5, %p58;
|
| 575 |
+
selp.b16 %rs58, 1, 0, %p71;
|
| 576 |
+
and.pred %p72, %p6, %p57;
|
| 577 |
+
selp.b16 %rs59, -1, 0, %p72;
|
| 578 |
+
shl.b16 %rs60, %rs59, 1;
|
| 579 |
+
or.b16 %rs61, %rs58, %rs60;
|
| 580 |
+
and.b16 %rs62, %rs61, 3;
|
| 581 |
+
or.b16 %rs63, %rs62, %rs57;
|
| 582 |
+
shl.b16 %rs64, %rs63, 4;
|
| 583 |
+
or.b16 %rs65, %rs52, %rs64;
|
| 584 |
+
.loc 1 32 62 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:62
|
| 585 |
+
and.b16 %rs66, %rs65, 255;
|
| 586 |
+
setp.eq.b16 %p73, %rs66, 0;
|
| 587 |
+
@%p73 bra $L__BB0_50;
|
| 588 |
+
bra.uni $L__BB0_49;
|
| 589 |
+
$L__BB0_50:
|
| 590 |
+
.loc 1 0 62 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0:62
|
| 591 |
+
ld.param.b64 %rd107, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_9];
|
| 592 |
+
ld.param.b64 %rd106, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_8];
|
| 593 |
+
ld.param.b64 %rd100, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_2];
|
| 594 |
+
.loc 1 24 19 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:24:19
|
| 595 |
+
rem.s64 %rd88, %rd1, %rd106;
|
| 596 |
+
rem.s64 %rd87, %rd2, %rd106;
|
| 597 |
+
rem.s64 %rd86, %rd3, %rd106;
|
| 598 |
+
rem.s64 %rd85, %rd4, %rd106;
|
| 599 |
+
rem.s64 %rd84, %rd5, %rd106;
|
| 600 |
+
rem.s64 %rd83, %rd6, %rd106;
|
| 601 |
+
rem.s64 %rd82, %rd7, %rd106;
|
| 602 |
+
rem.s64 %rd81, %rd8, %rd106;
|
| 603 |
+
.loc 1 32 62 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:62
|
| 604 |
+
bar.sync 0;
|
| 605 |
+
.loc 1 33 39 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:39
|
| 606 |
+
mul.lo.s64 %rd306, %rd73, %rd106;
|
| 607 |
+
mul.lo.s64 %rd307, %rd74, %rd106;
|
| 608 |
+
mul.lo.s64 %rd308, %rd75, %rd106;
|
| 609 |
+
mul.lo.s64 %rd309, %rd76, %rd106;
|
| 610 |
+
mul.lo.s64 %rd310, %rd77, %rd106;
|
| 611 |
+
mul.lo.s64 %rd311, %rd78, %rd106;
|
| 612 |
+
mul.lo.s64 %rd312, %rd79, %rd106;
|
| 613 |
+
mul.lo.s64 %rd313, %rd80, %rd106;
|
| 614 |
+
.loc 1 33 30 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:30
|
| 615 |
+
shl.b64 %rd314, %rd81, 1;
|
| 616 |
+
add.s64 %rd315, %rd100, %rd314;
|
| 617 |
+
shl.b64 %rd316, %rd306, 1;
|
| 618 |
+
add.s64 %rd235, %rd315, %rd316;
|
| 619 |
+
shl.b64 %rd317, %rd82, 1;
|
| 620 |
+
add.s64 %rd318, %rd100, %rd317;
|
| 621 |
+
shl.b64 %rd319, %rd307, 1;
|
| 622 |
+
add.s64 %rd238, %rd318, %rd319;
|
| 623 |
+
shl.b64 %rd320, %rd83, 1;
|
| 624 |
+
add.s64 %rd321, %rd100, %rd320;
|
| 625 |
+
shl.b64 %rd322, %rd308, 1;
|
| 626 |
+
add.s64 %rd241, %rd321, %rd322;
|
| 627 |
+
shl.b64 %rd323, %rd84, 1;
|
| 628 |
+
add.s64 %rd324, %rd100, %rd323;
|
| 629 |
+
shl.b64 %rd325, %rd309, 1;
|
| 630 |
+
add.s64 %rd244, %rd324, %rd325;
|
| 631 |
+
shl.b64 %rd326, %rd85, 1;
|
| 632 |
+
add.s64 %rd327, %rd100, %rd326;
|
| 633 |
+
shl.b64 %rd328, %rd310, 1;
|
| 634 |
+
add.s64 %rd247, %rd327, %rd328;
|
| 635 |
+
shl.b64 %rd329, %rd86, 1;
|
| 636 |
+
add.s64 %rd330, %rd100, %rd329;
|
| 637 |
+
shl.b64 %rd331, %rd311, 1;
|
| 638 |
+
add.s64 %rd250, %rd330, %rd331;
|
| 639 |
+
shl.b64 %rd332, %rd87, 1;
|
| 640 |
+
add.s64 %rd333, %rd100, %rd332;
|
| 641 |
+
shl.b64 %rd334, %rd312, 1;
|
| 642 |
+
add.s64 %rd253, %rd333, %rd334;
|
| 643 |
+
shl.b64 %rd335, %rd88, 1;
|
| 644 |
+
add.s64 %rd336, %rd100, %rd335;
|
| 645 |
+
shl.b64 %rd337, %rd313, 1;
|
| 646 |
+
add.s64 %rd256, %rd336, %rd337;
|
| 647 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 648 |
+
// begin inline asm
|
| 649 |
+
mov.u64 %rd234, 0x0;
|
| 650 |
+
createpolicy.fractional.L2::evict_last.b64 %rd234, 1.0;
|
| 651 |
+
// end inline asm
|
| 652 |
+
// begin inline asm
|
| 653 |
+
mov.u16 %rs67, 0x0;
|
| 654 |
+
@%p1 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs67 }, [ %rd235 + 0 ], %rd234;
|
| 655 |
+
// end inline asm
|
| 656 |
+
// begin inline asm
|
| 657 |
+
mov.u64 %rd237, 0x0;
|
| 658 |
+
createpolicy.fractional.L2::evict_last.b64 %rd237, 1.0;
|
| 659 |
+
// end inline asm
|
| 660 |
+
// begin inline asm
|
| 661 |
+
mov.u16 %rs68, 0x0;
|
| 662 |
+
@%p2 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs68 }, [ %rd238 + 0 ], %rd237;
|
| 663 |
+
// end inline asm
|
| 664 |
+
// begin inline asm
|
| 665 |
+
mov.u64 %rd240, 0x0;
|
| 666 |
+
createpolicy.fractional.L2::evict_last.b64 %rd240, 1.0;
|
| 667 |
+
// end inline asm
|
| 668 |
+
// begin inline asm
|
| 669 |
+
mov.u16 %rs69, 0x0;
|
| 670 |
+
@%p3 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs69 }, [ %rd241 + 0 ], %rd240;
|
| 671 |
+
// end inline asm
|
| 672 |
+
// begin inline asm
|
| 673 |
+
mov.u64 %rd243, 0x0;
|
| 674 |
+
createpolicy.fractional.L2::evict_last.b64 %rd243, 1.0;
|
| 675 |
+
// end inline asm
|
| 676 |
+
// begin inline asm
|
| 677 |
+
mov.u16 %rs70, 0x0;
|
| 678 |
+
@%p4 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs70 }, [ %rd244 + 0 ], %rd243;
|
| 679 |
+
// end inline asm
|
| 680 |
+
// begin inline asm
|
| 681 |
+
mov.u64 %rd246, 0x0;
|
| 682 |
+
createpolicy.fractional.L2::evict_last.b64 %rd246, 1.0;
|
| 683 |
+
// end inline asm
|
| 684 |
+
// begin inline asm
|
| 685 |
+
mov.u16 %rs71, 0x0;
|
| 686 |
+
@%p5 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs71 }, [ %rd247 + 0 ], %rd246;
|
| 687 |
+
// end inline asm
|
| 688 |
+
// begin inline asm
|
| 689 |
+
mov.u64 %rd249, 0x0;
|
| 690 |
+
createpolicy.fractional.L2::evict_last.b64 %rd249, 1.0;
|
| 691 |
+
// end inline asm
|
| 692 |
+
// begin inline asm
|
| 693 |
+
mov.u16 %rs72, 0x0;
|
| 694 |
+
@%p6 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs72 }, [ %rd250 + 0 ], %rd249;
|
| 695 |
+
// end inline asm
|
| 696 |
+
// begin inline asm
|
| 697 |
+
mov.u64 %rd252, 0x0;
|
| 698 |
+
createpolicy.fractional.L2::evict_last.b64 %rd252, 1.0;
|
| 699 |
+
// end inline asm
|
| 700 |
+
// begin inline asm
|
| 701 |
+
mov.u16 %rs73, 0x0;
|
| 702 |
+
@%p7 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs73 }, [ %rd253 + 0 ], %rd252;
|
| 703 |
+
// end inline asm
|
| 704 |
+
// begin inline asm
|
| 705 |
+
mov.u64 %rd255, 0x0;
|
| 706 |
+
createpolicy.fractional.L2::evict_last.b64 %rd255, 1.0;
|
| 707 |
+
// end inline asm
|
| 708 |
+
// begin inline asm
|
| 709 |
+
mov.u16 %rs74, 0x0;
|
| 710 |
+
@%p8 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs74 }, [ %rd256 + 0 ], %rd255;
|
| 711 |
+
// end inline asm
|
| 712 |
+
.loc 1 38 31 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:38:31
|
| 713 |
+
shr.u64 %rd338, %rd106, 63;
|
| 714 |
+
add.s64 %rd339, %rd106, %rd338;
|
| 715 |
+
shr.s64 %rd340, %rd339, 1;
|
| 716 |
+
.loc 1 38 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:38:18
|
| 717 |
+
sub.s64 %rd89, %rd106, %rd340;
|
| 718 |
+
.loc 1 39 19 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:39:19
|
| 719 |
+
setp.lt.s64 %p106, %rd81, %rd89;
|
| 720 |
+
setp.lt.s64 %p107, %rd82, %rd89;
|
| 721 |
+
setp.lt.s64 %p108, %rd83, %rd89;
|
| 722 |
+
setp.lt.s64 %p109, %rd84, %rd89;
|
| 723 |
+
setp.lt.s64 %p110, %rd85, %rd89;
|
| 724 |
+
setp.lt.s64 %p111, %rd86, %rd89;
|
| 725 |
+
setp.lt.s64 %p112, %rd87, %rd89;
|
| 726 |
+
setp.lt.s64 %p113, %rd88, %rd89;
|
| 727 |
+
.loc 1 40 35 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:35
|
| 728 |
+
sub.s64 %rd341, %rd8, %rd81;
|
| 729 |
+
sub.s64 %rd342, %rd7, %rd82;
|
| 730 |
+
sub.s64 %rd343, %rd6, %rd83;
|
| 731 |
+
sub.s64 %rd344, %rd5, %rd84;
|
| 732 |
+
sub.s64 %rd345, %rd4, %rd85;
|
| 733 |
+
sub.s64 %rd346, %rd3, %rd86;
|
| 734 |
+
sub.s64 %rd347, %rd2, %rd87;
|
| 735 |
+
sub.s64 %rd348, %rd1, %rd88;
|
| 736 |
+
.loc 1 40 31 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:31
|
| 737 |
+
shl.b64 %rd349, %rd341, 1;
|
| 738 |
+
add.s64 %rd350, %rd98, %rd349;
|
| 739 |
+
and.b64 %rd351, %rd339, -2;
|
| 740 |
+
add.s64 %rd352, %rd350, %rd351;
|
| 741 |
+
add.s64 %rd259, %rd352, %rd314;
|
| 742 |
+
shl.b64 %rd353, %rd342, 1;
|
| 743 |
+
add.s64 %rd354, %rd98, %rd353;
|
| 744 |
+
add.s64 %rd355, %rd354, %rd351;
|
| 745 |
+
add.s64 %rd262, %rd355, %rd317;
|
| 746 |
+
shl.b64 %rd356, %rd343, 1;
|
| 747 |
+
add.s64 %rd357, %rd98, %rd356;
|
| 748 |
+
add.s64 %rd358, %rd357, %rd351;
|
| 749 |
+
add.s64 %rd265, %rd358, %rd320;
|
| 750 |
+
shl.b64 %rd359, %rd344, 1;
|
| 751 |
+
add.s64 %rd360, %rd98, %rd359;
|
| 752 |
+
add.s64 %rd361, %rd360, %rd351;
|
| 753 |
+
add.s64 %rd268, %rd361, %rd323;
|
| 754 |
+
shl.b64 %rd362, %rd345, 1;
|
| 755 |
+
add.s64 %rd363, %rd98, %rd362;
|
| 756 |
+
add.s64 %rd364, %rd363, %rd351;
|
| 757 |
+
add.s64 %rd271, %rd364, %rd326;
|
| 758 |
+
shl.b64 %rd365, %rd346, 1;
|
| 759 |
+
add.s64 %rd366, %rd98, %rd365;
|
| 760 |
+
add.s64 %rd367, %rd366, %rd351;
|
| 761 |
+
add.s64 %rd274, %rd367, %rd329;
|
| 762 |
+
shl.b64 %rd368, %rd347, 1;
|
| 763 |
+
add.s64 %rd369, %rd98, %rd368;
|
| 764 |
+
add.s64 %rd370, %rd369, %rd351;
|
| 765 |
+
add.s64 %rd277, %rd370, %rd332;
|
| 766 |
+
shl.b64 %rd371, %rd348, 1;
|
| 767 |
+
add.s64 %rd372, %rd98, %rd371;
|
| 768 |
+
add.s64 %rd373, %rd372, %rd351;
|
| 769 |
+
add.s64 %rd280, %rd373, %rd335;
|
| 770 |
+
.loc 1 40 68 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:68
|
| 771 |
+
and.pred %p82, %p1, %p106;
|
| 772 |
+
and.pred %p83, %p2, %p107;
|
| 773 |
+
and.pred %p84, %p3, %p108;
|
| 774 |
+
and.pred %p85, %p4, %p109;
|
| 775 |
+
and.pred %p86, %p5, %p110;
|
| 776 |
+
and.pred %p87, %p6, %p111;
|
| 777 |
+
and.pred %p88, %p7, %p112;
|
| 778 |
+
and.pred %p89, %p8, %p113;
|
| 779 |
+
.loc 1 40 60 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:60
|
| 780 |
+
// begin inline asm
|
| 781 |
+
mov.u64 %rd258, 0x0;
|
| 782 |
+
createpolicy.fractional.L2::evict_last.b64 %rd258, 1.0;
|
| 783 |
+
// end inline asm
|
| 784 |
+
mov.b16 %rs76, 0;
|
| 785 |
+
// begin inline asm
|
| 786 |
+
mov.u16 %rs75, %rs76;
|
| 787 |
+
@%p82 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs75 }, [ %rd259 + 0 ], %rd258;
|
| 788 |
+
// end inline asm
|
| 789 |
+
// begin inline asm
|
| 790 |
+
mov.u64 %rd261, 0x0;
|
| 791 |
+
createpolicy.fractional.L2::evict_last.b64 %rd261, 1.0;
|
| 792 |
+
// end inline asm
|
| 793 |
+
// begin inline asm
|
| 794 |
+
mov.u16 %rs77, %rs76;
|
| 795 |
+
@%p83 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs77 }, [ %rd262 + 0 ], %rd261;
|
| 796 |
+
// end inline asm
|
| 797 |
+
// begin inline asm
|
| 798 |
+
mov.u64 %rd264, 0x0;
|
| 799 |
+
createpolicy.fractional.L2::evict_last.b64 %rd264, 1.0;
|
| 800 |
+
// end inline asm
|
| 801 |
+
// begin inline asm
|
| 802 |
+
mov.u16 %rs79, %rs76;
|
| 803 |
+
@%p84 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs79 }, [ %rd265 + 0 ], %rd264;
|
| 804 |
+
// end inline asm
|
| 805 |
+
// begin inline asm
|
| 806 |
+
mov.u64 %rd267, 0x0;
|
| 807 |
+
createpolicy.fractional.L2::evict_last.b64 %rd267, 1.0;
|
| 808 |
+
// end inline asm
|
| 809 |
+
// begin inline asm
|
| 810 |
+
mov.u16 %rs81, %rs76;
|
| 811 |
+
@%p85 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs81 }, [ %rd268 + 0 ], %rd267;
|
| 812 |
+
// end inline asm
|
| 813 |
+
// begin inline asm
|
| 814 |
+
mov.u64 %rd270, 0x0;
|
| 815 |
+
createpolicy.fractional.L2::evict_last.b64 %rd270, 1.0;
|
| 816 |
+
// end inline asm
|
| 817 |
+
// begin inline asm
|
| 818 |
+
mov.u16 %rs83, %rs76;
|
| 819 |
+
@%p86 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs83 }, [ %rd271 + 0 ], %rd270;
|
| 820 |
+
// end inline asm
|
| 821 |
+
// begin inline asm
|
| 822 |
+
mov.u64 %rd273, 0x0;
|
| 823 |
+
createpolicy.fractional.L2::evict_last.b64 %rd273, 1.0;
|
| 824 |
+
// end inline asm
|
| 825 |
+
// begin inline asm
|
| 826 |
+
mov.u16 %rs85, %rs76;
|
| 827 |
+
@%p87 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs85 }, [ %rd274 + 0 ], %rd273;
|
| 828 |
+
// end inline asm
|
| 829 |
+
// begin inline asm
|
| 830 |
+
mov.u64 %rd276, 0x0;
|
| 831 |
+
createpolicy.fractional.L2::evict_last.b64 %rd276, 1.0;
|
| 832 |
+
// end inline asm
|
| 833 |
+
// begin inline asm
|
| 834 |
+
mov.u16 %rs87, %rs76;
|
| 835 |
+
@%p88 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs87 }, [ %rd277 + 0 ], %rd276;
|
| 836 |
+
// end inline asm
|
| 837 |
+
// begin inline asm
|
| 838 |
+
mov.u64 %rd279, 0x0;
|
| 839 |
+
createpolicy.fractional.L2::evict_last.b64 %rd279, 1.0;
|
| 840 |
+
// end inline asm
|
| 841 |
+
// begin inline asm
|
| 842 |
+
mov.u16 %rs89, %rs76;
|
| 843 |
+
@%p89 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs89 }, [ %rd280 + 0 ], %rd279;
|
| 844 |
+
// end inline asm
|
| 845 |
+
.loc 1 44 20 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:44:20
|
| 846 |
+
setp.ge.s64 %p114, %rd88, %rd89;
|
| 847 |
+
setp.ge.s64 %p115, %rd87, %rd89;
|
| 848 |
+
setp.ge.s64 %p116, %rd86, %rd89;
|
| 849 |
+
setp.ge.s64 %p117, %rd85, %rd89;
|
| 850 |
+
setp.ge.s64 %p118, %rd84, %rd89;
|
| 851 |
+
setp.ge.s64 %p119, %rd83, %rd89;
|
| 852 |
+
setp.ge.s64 %p120, %rd82, %rd89;
|
| 853 |
+
setp.ge.s64 %p121, %rd81, %rd89;
|
| 854 |
+
.loc 1 47 47 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:47
|
| 855 |
+
sub.s64 %rd374, %rd81, %rd106;
|
| 856 |
+
sub.s64 %rd375, %rd82, %rd106;
|
| 857 |
+
sub.s64 %rd376, %rd83, %rd106;
|
| 858 |
+
sub.s64 %rd377, %rd84, %rd106;
|
| 859 |
+
sub.s64 %rd378, %rd85, %rd106;
|
| 860 |
+
sub.s64 %rd379, %rd86, %rd106;
|
| 861 |
+
sub.s64 %rd380, %rd87, %rd106;
|
| 862 |
+
sub.s64 %rd381, %rd88, %rd106;
|
| 863 |
+
.loc 1 47 31 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:31
|
| 864 |
+
shl.b64 %rd382, %rd374, 1;
|
| 865 |
+
add.s64 %rd283, %rd352, %rd382;
|
| 866 |
+
shl.b64 %rd383, %rd375, 1;
|
| 867 |
+
add.s64 %rd286, %rd355, %rd383;
|
| 868 |
+
shl.b64 %rd384, %rd376, 1;
|
| 869 |
+
add.s64 %rd289, %rd358, %rd384;
|
| 870 |
+
shl.b64 %rd385, %rd377, 1;
|
| 871 |
+
add.s64 %rd292, %rd361, %rd385;
|
| 872 |
+
shl.b64 %rd386, %rd378, 1;
|
| 873 |
+
add.s64 %rd295, %rd364, %rd386;
|
| 874 |
+
shl.b64 %rd387, %rd379, 1;
|
| 875 |
+
add.s64 %rd298, %rd367, %rd387;
|
| 876 |
+
shl.b64 %rd388, %rd380, 1;
|
| 877 |
+
add.s64 %rd301, %rd370, %rd388;
|
| 878 |
+
shl.b64 %rd389, %rd381, 1;
|
| 879 |
+
add.s64 %rd304, %rd373, %rd389;
|
| 880 |
+
.loc 1 47 81 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:81
|
| 881 |
+
and.pred %p90, %p1, %p121;
|
| 882 |
+
and.pred %p91, %p2, %p120;
|
| 883 |
+
and.pred %p92, %p3, %p119;
|
| 884 |
+
and.pred %p93, %p4, %p118;
|
| 885 |
+
and.pred %p94, %p5, %p117;
|
| 886 |
+
and.pred %p95, %p6, %p116;
|
| 887 |
+
and.pred %p96, %p7, %p115;
|
| 888 |
+
and.pred %p97, %p8, %p114;
|
| 889 |
+
.loc 1 47 73 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:73
|
| 890 |
+
// begin inline asm
|
| 891 |
+
mov.u64 %rd282, 0x0;
|
| 892 |
+
createpolicy.fractional.L2::evict_last.b64 %rd282, 1.0;
|
| 893 |
+
// end inline asm
|
| 894 |
+
// begin inline asm
|
| 895 |
+
mov.u16 %rs91, %rs76;
|
| 896 |
+
@%p90 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs91 }, [ %rd283 + 0 ], %rd282;
|
| 897 |
+
// end inline asm
|
| 898 |
+
// begin inline asm
|
| 899 |
+
mov.u64 %rd285, 0x0;
|
| 900 |
+
createpolicy.fractional.L2::evict_last.b64 %rd285, 1.0;
|
| 901 |
+
// end inline asm
|
| 902 |
+
// begin inline asm
|
| 903 |
+
mov.u16 %rs93, %rs76;
|
| 904 |
+
@%p91 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs93 }, [ %rd286 + 0 ], %rd285;
|
| 905 |
+
// end inline asm
|
| 906 |
+
// begin inline asm
|
| 907 |
+
mov.u64 %rd288, 0x0;
|
| 908 |
+
createpolicy.fractional.L2::evict_last.b64 %rd288, 1.0;
|
| 909 |
+
// end inline asm
|
| 910 |
+
// begin inline asm
|
| 911 |
+
mov.u16 %rs95, %rs76;
|
| 912 |
+
@%p92 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs95 }, [ %rd289 + 0 ], %rd288;
|
| 913 |
+
// end inline asm
|
| 914 |
+
// begin inline asm
|
| 915 |
+
mov.u64 %rd291, 0x0;
|
| 916 |
+
createpolicy.fractional.L2::evict_last.b64 %rd291, 1.0;
|
| 917 |
+
// end inline asm
|
| 918 |
+
// begin inline asm
|
| 919 |
+
mov.u16 %rs97, %rs76;
|
| 920 |
+
@%p93 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs97 }, [ %rd292 + 0 ], %rd291;
|
| 921 |
+
// end inline asm
|
| 922 |
+
// begin inline asm
|
| 923 |
+
mov.u64 %rd294, 0x0;
|
| 924 |
+
createpolicy.fractional.L2::evict_last.b64 %rd294, 1.0;
|
| 925 |
+
// end inline asm
|
| 926 |
+
// begin inline asm
|
| 927 |
+
mov.u16 %rs99, %rs76;
|
| 928 |
+
@%p94 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs99 }, [ %rd295 + 0 ], %rd294;
|
| 929 |
+
// end inline asm
|
| 930 |
+
// begin inline asm
|
| 931 |
+
mov.u64 %rd297, 0x0;
|
| 932 |
+
createpolicy.fractional.L2::evict_last.b64 %rd297, 1.0;
|
| 933 |
+
// end inline asm
|
| 934 |
+
// begin inline asm
|
| 935 |
+
mov.u16 %rs101, %rs76;
|
| 936 |
+
@%p95 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs101 }, [ %rd298 + 0 ], %rd297;
|
| 937 |
+
// end inline asm
|
| 938 |
+
// begin inline asm
|
| 939 |
+
mov.u64 %rd300, 0x0;
|
| 940 |
+
createpolicy.fractional.L2::evict_last.b64 %rd300, 1.0;
|
| 941 |
+
// end inline asm
|
| 942 |
+
// begin inline asm
|
| 943 |
+
mov.u16 %rs103, %rs76;
|
| 944 |
+
@%p96 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs103 }, [ %rd301 + 0 ], %rd300;
|
| 945 |
+
// end inline asm
|
| 946 |
+
// begin inline asm
|
| 947 |
+
mov.u64 %rd303, 0x0;
|
| 948 |
+
createpolicy.fractional.L2::evict_last.b64 %rd303, 1.0;
|
| 949 |
+
// end inline asm
|
| 950 |
+
// begin inline asm
|
| 951 |
+
mov.u16 %rs105, %rs76;
|
| 952 |
+
@%p97 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs105 }, [ %rd304 + 0 ], %rd303;
|
| 953 |
+
// end inline asm
|
| 954 |
+
.loc 1 51 34 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:51:34
|
| 955 |
+
and.b64 %rd391, %rd212, %rd107;
|
| 956 |
+
and.b64 %rd393, %rd214, %rd107;
|
| 957 |
+
and.b64 %rd395, %rd216, %rd107;
|
| 958 |
+
and.b64 %rd397, %rd218, %rd107;
|
| 959 |
+
and.b64 %rd399, %rd220, %rd107;
|
| 960 |
+
and.b64 %rd401, %rd222, %rd107;
|
| 961 |
+
and.b64 %rd403, %rd224, %rd107;
|
| 962 |
+
and.b64 %rd405, %rd226, %rd107;
|
| 963 |
+
add.s64 %rd95, %rd405, %rd185;
|
| 964 |
+
add.s64 %rd94, %rd403, %rd181;
|
| 965 |
+
add.s64 %rd97, %rd401, %rd193;
|
| 966 |
+
add.s64 %rd96, %rd399, %rd189;
|
| 967 |
+
add.s64 %rd91, %rd397, %rd169;
|
| 968 |
+
add.s64 %rd90, %rd395, %rd165;
|
| 969 |
+
add.s64 %rd93, %rd393, %rd177;
|
| 970 |
+
add.s64 %rd92, %rd391, %rd173;
|
| 971 |
+
.loc 1 52 28 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:28
|
| 972 |
+
setp.lt.s64 %p122, %rd92, 0;
|
| 973 |
+
setp.lt.s64 %p123, %rd93, 0;
|
| 974 |
+
setp.lt.s64 %p124, %rd90, 0;
|
| 975 |
+
setp.lt.s64 %p125, %rd91, 0;
|
| 976 |
+
setp.lt.s64 %p126, %rd96, 0;
|
| 977 |
+
setp.lt.s64 %p127, %rd97, 0;
|
| 978 |
+
setp.lt.s64 %p128, %rd94, 0;
|
| 979 |
+
setp.lt.s64 %p129, %rd95, 0;
|
| 980 |
+
.loc 1 52 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:46
|
| 981 |
+
setp.ge.s64 %p130, %rd92, %rd107;
|
| 982 |
+
setp.ge.s64 %p131, %rd93, %rd107;
|
| 983 |
+
setp.ge.s64 %p132, %rd90, %rd107;
|
| 984 |
+
setp.ge.s64 %p133, %rd91, %rd107;
|
| 985 |
+
setp.ge.s64 %p134, %rd96, %rd107;
|
| 986 |
+
setp.ge.s64 %p135, %rd97, %rd107;
|
| 987 |
+
setp.ge.s64 %p136, %rd94, %rd107;
|
| 988 |
+
setp.ge.s64 %p137, %rd95, %rd107;
|
| 989 |
+
.loc 1 52 38 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:38
|
| 990 |
+
or.pred %p138, %p129, %p137;
|
| 991 |
+
or.pred %p139, %p128, %p136;
|
| 992 |
+
or.pred %p140, %p127, %p135;
|
| 993 |
+
or.pred %p141, %p126, %p134;
|
| 994 |
+
or.pred %p142, %p125, %p133;
|
| 995 |
+
or.pred %p143, %p124, %p132;
|
| 996 |
+
or.pred %p144, %p123, %p131;
|
| 997 |
+
or.pred %p145, %p122, %p130;
|
| 998 |
+
.loc 1 52 54 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:54
|
| 999 |
+
and.pred %p146, %p3, %p145;
|
| 1000 |
+
selp.b16 %rs107, 1, 0, %p146;
|
| 1001 |
+
shl.b16 %rs108, %rs107, 2;
|
| 1002 |
+
and.pred %p147, %p4, %p144;
|
| 1003 |
+
selp.b16 %rs109, -1, 0, %p147;
|
| 1004 |
+
shl.b16 %rs110, %rs109, 3;
|
| 1005 |
+
or.b16 %rs111, %rs110, %rs108;
|
| 1006 |
+
and.pred %p148, %p1, %p143;
|
| 1007 |
+
selp.b16 %rs112, 1, 0, %p148;
|
| 1008 |
+
and.pred %p149, %p2, %p142;
|
| 1009 |
+
selp.b16 %rs113, -1, 0, %p149;
|
| 1010 |
+
shl.b16 %rs114, %rs113, 1;
|
| 1011 |
+
or.b16 %rs115, %rs112, %rs114;
|
| 1012 |
+
and.b16 %rs116, %rs115, 3;
|
| 1013 |
+
or.b16 %rs117, %rs116, %rs111;
|
| 1014 |
+
and.b16 %rs118, %rs117, 15;
|
| 1015 |
+
and.pred %p150, %p7, %p141;
|
| 1016 |
+
selp.b16 %rs119, 1, 0, %p150;
|
| 1017 |
+
shl.b16 %rs120, %rs119, 2;
|
| 1018 |
+
and.pred %p151, %p8, %p140;
|
| 1019 |
+
selp.b16 %rs121, -1, 0, %p151;
|
| 1020 |
+
shl.b16 %rs122, %rs121, 3;
|
| 1021 |
+
or.b16 %rs123, %rs122, %rs120;
|
| 1022 |
+
and.pred %p152, %p5, %p139;
|
| 1023 |
+
selp.b16 %rs124, 1, 0, %p152;
|
| 1024 |
+
and.pred %p153, %p6, %p138;
|
| 1025 |
+
selp.b16 %rs125, -1, 0, %p153;
|
| 1026 |
+
shl.b16 %rs126, %rs125, 1;
|
| 1027 |
+
or.b16 %rs127, %rs124, %rs126;
|
| 1028 |
+
and.b16 %rs128, %rs127, 3;
|
| 1029 |
+
or.b16 %rs129, %rs128, %rs123;
|
| 1030 |
+
shl.b16 %rs130, %rs129, 4;
|
| 1031 |
+
or.b16 %rs131, %rs118, %rs130;
|
| 1032 |
+
.loc 1 52 64 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:64
|
| 1033 |
+
and.b16 %rs132, %rs131, 255;
|
| 1034 |
+
setp.eq.b16 %p154, %rs132, 0;
|
| 1035 |
+
@%p154 bra $L__BB0_52;
|
| 1036 |
+
bra.uni $L__BB0_51;
|
| 1037 |
+
$L__BB0_52:
|
| 1038 |
+
.loc 1 0 64 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0:64
|
| 1039 |
+
ld.param.b64 %rd102, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_4];
|
| 1040 |
+
ld.param.b64 %rd101, [triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0_param_3];
|
| 1041 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1042 |
+
cvt.f32.bf16 %r79, %rs89;
|
| 1043 |
+
mov.b32 %r80, 0f00000000;
|
| 1044 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1045 |
+
sub.f32 %r81, %r80, %r79;
|
| 1046 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1047 |
+
cvt.f32.bf16 %r82, %rs105;
|
| 1048 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1049 |
+
selp.f32 %r83, %r81, %r82, %p113;
|
| 1050 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1051 |
+
cvt.f32.bf16 %r84, %rs87;
|
| 1052 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1053 |
+
sub.f32 %r85, %r80, %r84;
|
| 1054 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1055 |
+
cvt.f32.bf16 %r86, %rs103;
|
| 1056 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1057 |
+
selp.f32 %r87, %r85, %r86, %p112;
|
| 1058 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1059 |
+
cvt.f32.bf16 %r88, %rs85;
|
| 1060 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1061 |
+
sub.f32 %r89, %r80, %r88;
|
| 1062 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1063 |
+
cvt.f32.bf16 %r90, %rs101;
|
| 1064 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1065 |
+
selp.f32 %r91, %r89, %r90, %p111;
|
| 1066 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1067 |
+
cvt.f32.bf16 %r92, %rs83;
|
| 1068 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1069 |
+
sub.f32 %r93, %r80, %r92;
|
| 1070 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1071 |
+
cvt.f32.bf16 %r94, %rs99;
|
| 1072 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1073 |
+
selp.f32 %r95, %r93, %r94, %p110;
|
| 1074 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1075 |
+
cvt.f32.bf16 %r96, %rs81;
|
| 1076 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1077 |
+
sub.f32 %r97, %r80, %r96;
|
| 1078 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1079 |
+
cvt.f32.bf16 %r98, %rs97;
|
| 1080 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1081 |
+
selp.f32 %r99, %r97, %r98, %p109;
|
| 1082 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1083 |
+
cvt.f32.bf16 %r100, %rs79;
|
| 1084 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1085 |
+
sub.f32 %r101, %r80, %r100;
|
| 1086 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1087 |
+
cvt.f32.bf16 %r102, %rs95;
|
| 1088 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1089 |
+
selp.f32 %r103, %r101, %r102, %p108;
|
| 1090 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1091 |
+
cvt.f32.bf16 %r104, %rs77;
|
| 1092 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1093 |
+
sub.f32 %r105, %r80, %r104;
|
| 1094 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1095 |
+
cvt.f32.bf16 %r106, %rs93;
|
| 1096 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1097 |
+
selp.f32 %r107, %r105, %r106, %p107;
|
| 1098 |
+
.loc 1 40 119 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:40:119
|
| 1099 |
+
cvt.f32.bf16 %r108, %rs75;
|
| 1100 |
+
.loc 1 41 13 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:41:13
|
| 1101 |
+
sub.f32 %r109, %r80, %r108;
|
| 1102 |
+
.loc 1 47 132 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:47:132
|
| 1103 |
+
cvt.f32.bf16 %r110, %rs91;
|
| 1104 |
+
.loc 1 0 0 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:0
|
| 1105 |
+
selp.f32 %r111, %r109, %r110, %p106;
|
| 1106 |
+
.loc 1 26 75 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:26:75
|
| 1107 |
+
cvt.f32.bf16 %r112, %rs40;
|
| 1108 |
+
cvt.f32.bf16 %r113, %rs39;
|
| 1109 |
+
cvt.f32.bf16 %r114, %rs38;
|
| 1110 |
+
cvt.f32.bf16 %r115, %rs37;
|
| 1111 |
+
cvt.f32.bf16 %r116, %rs36;
|
| 1112 |
+
cvt.f32.bf16 %r117, %rs35;
|
| 1113 |
+
cvt.f32.bf16 %r118, %rs34;
|
| 1114 |
+
cvt.f32.bf16 %r119, %rs33;
|
| 1115 |
+
.loc 1 52 64 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:64
|
| 1116 |
+
bar.sync 0;
|
| 1117 |
+
.loc 1 53 40 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:53:40
|
| 1118 |
+
mul.lo.s64 %rd444, %rd90, %rd106;
|
| 1119 |
+
mul.lo.s64 %rd445, %rd91, %rd106;
|
| 1120 |
+
mul.lo.s64 %rd446, %rd92, %rd106;
|
| 1121 |
+
mul.lo.s64 %rd447, %rd93, %rd106;
|
| 1122 |
+
mul.lo.s64 %rd448, %rd94, %rd106;
|
| 1123 |
+
mul.lo.s64 %rd449, %rd95, %rd106;
|
| 1124 |
+
mul.lo.s64 %rd450, %rd96, %rd106;
|
| 1125 |
+
mul.lo.s64 %rd451, %rd97, %rd106;
|
| 1126 |
+
.loc 1 53 31 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:53:31
|
| 1127 |
+
add.s64 %rd453, %rd101, %rd314;
|
| 1128 |
+
shl.b64 %rd454, %rd444, 1;
|
| 1129 |
+
add.s64 %rd413, %rd453, %rd454;
|
| 1130 |
+
add.s64 %rd456, %rd101, %rd317;
|
| 1131 |
+
shl.b64 %rd457, %rd445, 1;
|
| 1132 |
+
add.s64 %rd416, %rd456, %rd457;
|
| 1133 |
+
add.s64 %rd459, %rd101, %rd320;
|
| 1134 |
+
shl.b64 %rd460, %rd446, 1;
|
| 1135 |
+
add.s64 %rd419, %rd459, %rd460;
|
| 1136 |
+
add.s64 %rd462, %rd101, %rd323;
|
| 1137 |
+
shl.b64 %rd463, %rd447, 1;
|
| 1138 |
+
add.s64 %rd422, %rd462, %rd463;
|
| 1139 |
+
add.s64 %rd465, %rd101, %rd326;
|
| 1140 |
+
shl.b64 %rd466, %rd448, 1;
|
| 1141 |
+
add.s64 %rd425, %rd465, %rd466;
|
| 1142 |
+
add.s64 %rd468, %rd101, %rd329;
|
| 1143 |
+
shl.b64 %rd469, %rd449, 1;
|
| 1144 |
+
add.s64 %rd428, %rd468, %rd469;
|
| 1145 |
+
add.s64 %rd471, %rd101, %rd332;
|
| 1146 |
+
shl.b64 %rd472, %rd450, 1;
|
| 1147 |
+
add.s64 %rd431, %rd471, %rd472;
|
| 1148 |
+
add.s64 %rd474, %rd101, %rd335;
|
| 1149 |
+
shl.b64 %rd475, %rd451, 1;
|
| 1150 |
+
add.s64 %rd434, %rd474, %rd475;
|
| 1151 |
+
.loc 1 53 48 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:53:48
|
| 1152 |
+
// begin inline asm
|
| 1153 |
+
mov.u64 %rd414, 0x0;
|
| 1154 |
+
createpolicy.fractional.L2::evict_last.b64 %rd414, 1.0;
|
| 1155 |
+
// end inline asm
|
| 1156 |
+
// begin inline asm
|
| 1157 |
+
mov.u16 %rs133, 0x0;
|
| 1158 |
+
@%p1 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs133 }, [ %rd413 + 0 ], %rd414;
|
| 1159 |
+
// end inline asm
|
| 1160 |
+
// begin inline asm
|
| 1161 |
+
mov.u64 %rd417, 0x0;
|
| 1162 |
+
createpolicy.fractional.L2::evict_last.b64 %rd417, 1.0;
|
| 1163 |
+
// end inline asm
|
| 1164 |
+
// begin inline asm
|
| 1165 |
+
mov.u16 %rs134, 0x0;
|
| 1166 |
+
@%p2 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs134 }, [ %rd416 + 0 ], %rd417;
|
| 1167 |
+
// end inline asm
|
| 1168 |
+
// begin inline asm
|
| 1169 |
+
mov.u64 %rd420, 0x0;
|
| 1170 |
+
createpolicy.fractional.L2::evict_last.b64 %rd420, 1.0;
|
| 1171 |
+
// end inline asm
|
| 1172 |
+
// begin inline asm
|
| 1173 |
+
mov.u16 %rs135, 0x0;
|
| 1174 |
+
@%p3 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs135 }, [ %rd419 + 0 ], %rd420;
|
| 1175 |
+
// end inline asm
|
| 1176 |
+
// begin inline asm
|
| 1177 |
+
mov.u64 %rd423, 0x0;
|
| 1178 |
+
createpolicy.fractional.L2::evict_last.b64 %rd423, 1.0;
|
| 1179 |
+
// end inline asm
|
| 1180 |
+
// begin inline asm
|
| 1181 |
+
mov.u16 %rs136, 0x0;
|
| 1182 |
+
@%p4 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs136 }, [ %rd422 + 0 ], %rd423;
|
| 1183 |
+
// end inline asm
|
| 1184 |
+
// begin inline asm
|
| 1185 |
+
mov.u64 %rd426, 0x0;
|
| 1186 |
+
createpolicy.fractional.L2::evict_last.b64 %rd426, 1.0;
|
| 1187 |
+
// end inline asm
|
| 1188 |
+
// begin inline asm
|
| 1189 |
+
mov.u16 %rs137, 0x0;
|
| 1190 |
+
@%p5 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs137 }, [ %rd425 + 0 ], %rd426;
|
| 1191 |
+
// end inline asm
|
| 1192 |
+
// begin inline asm
|
| 1193 |
+
mov.u64 %rd429, 0x0;
|
| 1194 |
+
createpolicy.fractional.L2::evict_last.b64 %rd429, 1.0;
|
| 1195 |
+
// end inline asm
|
| 1196 |
+
// begin inline asm
|
| 1197 |
+
mov.u16 %rs138, 0x0;
|
| 1198 |
+
@%p6 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs138 }, [ %rd428 + 0 ], %rd429;
|
| 1199 |
+
// end inline asm
|
| 1200 |
+
// begin inline asm
|
| 1201 |
+
mov.u64 %rd432, 0x0;
|
| 1202 |
+
createpolicy.fractional.L2::evict_last.b64 %rd432, 1.0;
|
| 1203 |
+
// end inline asm
|
| 1204 |
+
// begin inline asm
|
| 1205 |
+
mov.u16 %rs139, 0x0;
|
| 1206 |
+
@%p7 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs139 }, [ %rd431 + 0 ], %rd432;
|
| 1207 |
+
// end inline asm
|
| 1208 |
+
// begin inline asm
|
| 1209 |
+
mov.u64 %rd435, 0x0;
|
| 1210 |
+
createpolicy.fractional.L2::evict_last.b64 %rd435, 1.0;
|
| 1211 |
+
// end inline asm
|
| 1212 |
+
// begin inline asm
|
| 1213 |
+
mov.u16 %rs140, 0x0;
|
| 1214 |
+
@%p8 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs140 }, [ %rd434 + 0 ], %rd435;
|
| 1215 |
+
// end inline asm
|
| 1216 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1217 |
+
mov.b32 %r120, {%rs67, %rs133};
|
| 1218 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1219 |
+
mov.b32 {%rs149, %rs150}, %r120;
|
| 1220 |
+
cvt.f32.bf16 %r121, %rs149;
|
| 1221 |
+
cvt.f32.bf16 %r122, %rs150;
|
| 1222 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1223 |
+
mul.f32 %r123, %r111, %r122;
|
| 1224 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1225 |
+
mov.b32 %r124, {%rs68, %rs134};
|
| 1226 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1227 |
+
mov.b32 {%rs151, %rs152}, %r124;
|
| 1228 |
+
cvt.f32.bf16 %r125, %rs151;
|
| 1229 |
+
cvt.f32.bf16 %r126, %rs152;
|
| 1230 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1231 |
+
mul.f32 %r127, %r107, %r126;
|
| 1232 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1233 |
+
mov.b32 %r128, {%rs69, %rs135};
|
| 1234 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1235 |
+
mov.b32 {%rs153, %rs154}, %r128;
|
| 1236 |
+
cvt.f32.bf16 %r129, %rs153;
|
| 1237 |
+
cvt.f32.bf16 %r130, %rs154;
|
| 1238 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1239 |
+
mul.f32 %r131, %r103, %r130;
|
| 1240 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1241 |
+
mov.b32 %r132, {%rs70, %rs136};
|
| 1242 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1243 |
+
mov.b32 {%rs155, %rs156}, %r132;
|
| 1244 |
+
cvt.f32.bf16 %r133, %rs155;
|
| 1245 |
+
cvt.f32.bf16 %r134, %rs156;
|
| 1246 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1247 |
+
mul.f32 %r135, %r99, %r134;
|
| 1248 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1249 |
+
mov.b32 %r136, {%rs71, %rs137};
|
| 1250 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1251 |
+
mov.b32 {%rs157, %rs158}, %r136;
|
| 1252 |
+
cvt.f32.bf16 %r137, %rs157;
|
| 1253 |
+
cvt.f32.bf16 %r138, %rs158;
|
| 1254 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1255 |
+
mul.f32 %r139, %r95, %r138;
|
| 1256 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1257 |
+
mov.b32 %r140, {%rs72, %rs138};
|
| 1258 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1259 |
+
mov.b32 {%rs159, %rs160}, %r140;
|
| 1260 |
+
cvt.f32.bf16 %r141, %rs159;
|
| 1261 |
+
cvt.f32.bf16 %r142, %rs160;
|
| 1262 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1263 |
+
mul.f32 %r143, %r91, %r142;
|
| 1264 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1265 |
+
mov.b32 %r144, {%rs73, %rs139};
|
| 1266 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1267 |
+
mov.b32 {%rs161, %rs162}, %r144;
|
| 1268 |
+
cvt.f32.bf16 %r145, %rs161;
|
| 1269 |
+
cvt.f32.bf16 %r146, %rs162;
|
| 1270 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1271 |
+
mul.f32 %r147, %r87, %r146;
|
| 1272 |
+
.loc 1 33 46 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:46
|
| 1273 |
+
mov.b32 %r148, {%rs74, %rs140};
|
| 1274 |
+
.loc 1 33 86 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:33:86
|
| 1275 |
+
mov.b32 {%rs163, %rs164}, %r148;
|
| 1276 |
+
cvt.f32.bf16 %r149, %rs163;
|
| 1277 |
+
cvt.f32.bf16 %r150, %rs164;
|
| 1278 |
+
.loc 1 34 18 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:34:18
|
| 1279 |
+
mul.f32 %r151, %r83, %r150;
|
| 1280 |
+
.loc 1 55 19 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:55:19
|
| 1281 |
+
fma.rn.f32 %r152, %r119, %r121, %r123;
|
| 1282 |
+
fma.rn.f32 %r153, %r118, %r125, %r127;
|
| 1283 |
+
fma.rn.f32 %r154, %r117, %r129, %r131;
|
| 1284 |
+
fma.rn.f32 %r155, %r116, %r133, %r135;
|
| 1285 |
+
fma.rn.f32 %r156, %r115, %r137, %r139;
|
| 1286 |
+
fma.rn.f32 %r157, %r114, %r141, %r143;
|
| 1287 |
+
fma.rn.f32 %r158, %r113, %r145, %r147;
|
| 1288 |
+
fma.rn.f32 %r159, %r112, %r149, %r151;
|
| 1289 |
+
.loc 1 56 25 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:56:25
|
| 1290 |
+
add.s64 %rd436, %rd102, %rd196;
|
| 1291 |
+
add.s64 %rd437, %rd102, %rd197;
|
| 1292 |
+
add.s64 %rd438, %rd102, %rd198;
|
| 1293 |
+
add.s64 %rd439, %rd102, %rd199;
|
| 1294 |
+
add.s64 %rd440, %rd102, %rd200;
|
| 1295 |
+
add.s64 %rd441, %rd102, %rd201;
|
| 1296 |
+
add.s64 %rd442, %rd102, %rd202;
|
| 1297 |
+
add.s64 %rd443, %rd102, %rd203;
|
| 1298 |
+
.loc 1 56 37 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:56:37
|
| 1299 |
+
cvt.rn.bf16.f32 %rs141, %r152;
|
| 1300 |
+
cvt.rn.bf16.f32 %rs142, %r153;
|
| 1301 |
+
cvt.rn.bf16.f32 %rs143, %r154;
|
| 1302 |
+
cvt.rn.bf16.f32 %rs144, %r155;
|
| 1303 |
+
cvt.rn.bf16.f32 %rs145, %r156;
|
| 1304 |
+
cvt.rn.bf16.f32 %rs146, %r157;
|
| 1305 |
+
cvt.rn.bf16.f32 %rs147, %r158;
|
| 1306 |
+
cvt.rn.bf16.f32 %rs148, %r159;
|
| 1307 |
+
// begin inline asm
|
| 1308 |
+
@%p1 st.global.b16 [ %rd436 + 0 ], { %rs141 };
|
| 1309 |
+
// end inline asm
|
| 1310 |
+
// begin inline asm
|
| 1311 |
+
@%p2 st.global.b16 [ %rd437 + 0 ], { %rs142 };
|
| 1312 |
+
// end inline asm
|
| 1313 |
+
// begin inline asm
|
| 1314 |
+
@%p3 st.global.b16 [ %rd438 + 0 ], { %rs143 };
|
| 1315 |
+
// end inline asm
|
| 1316 |
+
// begin inline asm
|
| 1317 |
+
@%p4 st.global.b16 [ %rd439 + 0 ], { %rs144 };
|
| 1318 |
+
// end inline asm
|
| 1319 |
+
// begin inline asm
|
| 1320 |
+
@%p5 st.global.b16 [ %rd440 + 0 ], { %rs145 };
|
| 1321 |
+
// end inline asm
|
| 1322 |
+
// begin inline asm
|
| 1323 |
+
@%p6 st.global.b16 [ %rd441 + 0 ], { %rs146 };
|
| 1324 |
+
// end inline asm
|
| 1325 |
+
// begin inline asm
|
| 1326 |
+
@%p7 st.global.b16 [ %rd442 + 0 ], { %rs147 };
|
| 1327 |
+
// end inline asm
|
| 1328 |
+
// begin inline asm
|
| 1329 |
+
@%p8 st.global.b16 [ %rd443 + 0 ], { %rs148 };
|
| 1330 |
+
// end inline asm
|
| 1331 |
+
.loc 1 56 4 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:56:4
|
| 1332 |
+
ret;
|
| 1333 |
+
$L__BB0_49:
|
| 1334 |
+
.loc 1 32 62 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:32:62
|
| 1335 |
+
{ // callseq 0, 0
|
| 1336 |
+
.param .b64 param0;
|
| 1337 |
+
.param .b64 param1;
|
| 1338 |
+
.param .b32 param2;
|
| 1339 |
+
.param .b64 param3;
|
| 1340 |
+
.param .b64 param4;
|
| 1341 |
+
mov.b64 %rd228, assertFunc_0;
|
| 1342 |
+
cvta.global.u64 %rd229, %rd228;
|
| 1343 |
+
st.param.b64 [param3], %rd229;
|
| 1344 |
+
mov.b64 %rd230, assertFile_0;
|
| 1345 |
+
cvta.global.u64 %rd231, %rd230;
|
| 1346 |
+
st.param.b64 [param1], %rd231;
|
| 1347 |
+
mov.b64 %rd232, assertMessage_0;
|
| 1348 |
+
cvta.global.u64 %rd233, %rd232;
|
| 1349 |
+
st.param.b64 [param0], %rd233;
|
| 1350 |
+
st.param.b64 [param4], 1;
|
| 1351 |
+
st.param.b32 [param2], 32;
|
| 1352 |
+
call.uni __assertfail, (param0, param1, param2, param3, param4);
|
| 1353 |
+
} // callseq 0
|
| 1354 |
+
trap;
|
| 1355 |
+
$L__BB0_51:
|
| 1356 |
+
.loc 1 52 64 // crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py:52:64
|
| 1357 |
+
{ // callseq 1, 0
|
| 1358 |
+
.param .b64 param0;
|
| 1359 |
+
.param .b64 param1;
|
| 1360 |
+
.param .b32 param2;
|
| 1361 |
+
.param .b64 param3;
|
| 1362 |
+
.param .b64 param4;
|
| 1363 |
+
mov.b64 %rd406, assertFunc_1;
|
| 1364 |
+
cvta.global.u64 %rd407, %rd406;
|
| 1365 |
+
st.param.b64 [param3], %rd407;
|
| 1366 |
+
mov.b64 %rd408, assertFile_1;
|
| 1367 |
+
cvta.global.u64 %rd409, %rd408;
|
| 1368 |
+
st.param.b64 [param1], %rd409;
|
| 1369 |
+
mov.b64 %rd410, assertMessage_1;
|
| 1370 |
+
cvta.global.u64 %rd411, %rd410;
|
| 1371 |
+
st.param.b64 [param0], %rd411;
|
| 1372 |
+
st.param.b64 [param4], 1;
|
| 1373 |
+
st.param.b32 [param2], 52;
|
| 1374 |
+
call.uni __assertfail, (param0, param1, param2, param3, param4);
|
| 1375 |
+
} // callseq 1
|
| 1376 |
+
trap;
|
| 1377 |
+
$L__tmp1:
|
| 1378 |
+
$L__func_end0:
|
| 1379 |
+
// -- End function
|
| 1380 |
+
}
|
| 1381 |
+
.file 1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py"
|
| 1382 |
+
.section .debug_abbrev
|
| 1383 |
+
{
|
| 1384 |
+
.b8 1 // Abbreviation Code
|
| 1385 |
+
.b8 17 // DW_TAG_compile_unit
|
| 1386 |
+
.b8 0 // DW_CHILDREN_no
|
| 1387 |
+
.b8 37 // DW_AT_producer
|
| 1388 |
+
.b8 8 // DW_FORM_string
|
| 1389 |
+
.b8 19 // DW_AT_language
|
| 1390 |
+
.b8 5 // DW_FORM_data2
|
| 1391 |
+
.b8 3 // DW_AT_name
|
| 1392 |
+
.b8 8 // DW_FORM_string
|
| 1393 |
+
.b8 16 // DW_AT_stmt_list
|
| 1394 |
+
.b8 6 // DW_FORM_data4
|
| 1395 |
+
.b8 27 // DW_AT_comp_dir
|
| 1396 |
+
.b8 8 // DW_FORM_string
|
| 1397 |
+
.b8 0 // EOM(1)
|
| 1398 |
+
.b8 0 // EOM(2)
|
| 1399 |
+
.b8 0 // EOM(3)
|
| 1400 |
+
}
|
| 1401 |
+
.section .debug_info
|
| 1402 |
+
{
|
| 1403 |
+
.b32 135 // Length of Unit
|
| 1404 |
+
.b8 2 // DWARF version number
|
| 1405 |
+
.b8 0
|
| 1406 |
+
.b32 .debug_abbrev // Offset Into Abbrev. Section
|
| 1407 |
+
.b8 8 // Address Size (in bytes)
|
| 1408 |
+
.b8 1 // Abbrev [1] 0xb:0x80 DW_TAG_compile_unit
|
| 1409 |
+
.b8 116 // DW_AT_producer
|
| 1410 |
+
.b8 114
|
| 1411 |
+
.b8 105
|
| 1412 |
+
.b8 116
|
| 1413 |
+
.b8 111
|
| 1414 |
+
.b8 110
|
| 1415 |
+
.b8 0
|
| 1416 |
+
.b8 2 // DW_AT_language
|
| 1417 |
+
.b8 0
|
| 1418 |
+
.b8 99 // DW_AT_name
|
| 1419 |
+
.b8 114
|
| 1420 |
+
.b8 118
|
| 1421 |
+
.b8 120
|
| 1422 |
+
.b8 98
|
| 1423 |
+
.b8 122
|
| 1424 |
+
.b8 116
|
| 1425 |
+
.b8 107
|
| 1426 |
+
.b8 114
|
| 1427 |
+
.b8 51
|
| 1428 |
+
.b8 55
|
| 1429 |
+
.b8 50
|
| 1430 |
+
.b8 102
|
| 1431 |
+
.b8 103
|
| 1432 |
+
.b8 100
|
| 1433 |
+
.b8 110
|
| 1434 |
+
.b8 55
|
| 1435 |
+
.b8 98
|
| 1436 |
+
.b8 103
|
| 1437 |
+
.b8 114
|
| 1438 |
+
.b8 117
|
| 1439 |
+
.b8 100
|
| 1440 |
+
.b8 50
|
| 1441 |
+
.b8 50
|
| 1442 |
+
.b8 115
|
| 1443 |
+
.b8 51
|
| 1444 |
+
.b8 119
|
| 1445 |
+
.b8 109
|
| 1446 |
+
.b8 100
|
| 1447 |
+
.b8 50
|
| 1448 |
+
.b8 105
|
| 1449 |
+
.b8 115
|
| 1450 |
+
.b8 105
|
| 1451 |
+
.b8 100
|
| 1452 |
+
.b8 119
|
| 1453 |
+
.b8 111
|
| 1454 |
+
.b8 52
|
| 1455 |
+
.b8 101
|
| 1456 |
+
.b8 107
|
| 1457 |
+
.b8 52
|
| 1458 |
+
.b8 104
|
| 1459 |
+
.b8 108
|
| 1460 |
+
.b8 100
|
| 1461 |
+
.b8 110
|
| 1462 |
+
.b8 53
|
| 1463 |
+
.b8 54
|
| 1464 |
+
.b8 116
|
| 1465 |
+
.b8 117
|
| 1466 |
+
.b8 118
|
| 1467 |
+
.b8 50
|
| 1468 |
+
.b8 100
|
| 1469 |
+
.b8 106
|
| 1470 |
+
.b8 46
|
| 1471 |
+
.b8 112
|
| 1472 |
+
.b8 121
|
| 1473 |
+
.b8 0
|
| 1474 |
+
.b32 .debug_line // DW_AT_stmt_list
|
| 1475 |
+
.b8 47 // DW_AT_comp_dir
|
| 1476 |
+
.b8 119
|
| 1477 |
+
.b8 111
|
| 1478 |
+
.b8 114
|
| 1479 |
+
.b8 107
|
| 1480 |
+
.b8 115
|
| 1481 |
+
.b8 112
|
| 1482 |
+
.b8 97
|
| 1483 |
+
.b8 99
|
| 1484 |
+
.b8 101
|
| 1485 |
+
.b8 47
|
| 1486 |
+
.b8 104
|
| 1487 |
+
.b8 97
|
| 1488 |
+
.b8 110
|
| 1489 |
+
.b8 114
|
| 1490 |
+
.b8 117
|
| 1491 |
+
.b8 105
|
| 1492 |
+
.b8 47
|
| 1493 |
+
.b8 83
|
| 1494 |
+
.b8 112
|
| 1495 |
+
.b8 101
|
| 1496 |
+
.b8 99
|
| 1497 |
+
.b8 70
|
| 1498 |
+
.b8 111
|
| 1499 |
+
.b8 114
|
| 1500 |
+
.b8 103
|
| 1501 |
+
.b8 101
|
| 1502 |
+
.b8 45
|
| 1503 |
+
.b8 101
|
| 1504 |
+
.b8 120
|
| 1505 |
+
.b8 116
|
| 1506 |
+
.b8 47
|
| 1507 |
+
.b8 99
|
| 1508 |
+
.b8 97
|
| 1509 |
+
.b8 99
|
| 1510 |
+
.b8 104
|
| 1511 |
+
.b8 101
|
| 1512 |
+
.b8 47
|
| 1513 |
+
.b8 99
|
| 1514 |
+
.b8 111
|
| 1515 |
+
.b8 109
|
| 1516 |
+
.b8 112
|
| 1517 |
+
.b8 105
|
| 1518 |
+
.b8 108
|
| 1519 |
+
.b8 101
|
| 1520 |
+
.b8 100
|
| 1521 |
+
.b8 95
|
| 1522 |
+
.b8 107
|
| 1523 |
+
.b8 101
|
| 1524 |
+
.b8 114
|
| 1525 |
+
.b8 110
|
| 1526 |
+
.b8 101
|
| 1527 |
+
.b8 108
|
| 1528 |
+
.b8 115
|
| 1529 |
+
.b8 47
|
| 1530 |
+
.b8 114
|
| 1531 |
+
.b8 118
|
| 1532 |
+
.b8 0
|
| 1533 |
+
}
|
| 1534 |
+
.section .debug_macinfo { }
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.source
ADDED
|
@@ -0,0 +1,299 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":18:0)
|
| 2 |
+
#loc78 = loc("in_ptr0"(#loc))
|
| 3 |
+
#loc79 = loc("in_ptr1"(#loc))
|
| 4 |
+
#loc80 = loc("in_ptr2"(#loc))
|
| 5 |
+
#loc81 = loc("in_ptr3"(#loc))
|
| 6 |
+
#loc82 = loc("out_ptr0"(#loc))
|
| 7 |
+
#loc83 = loc("ks0"(#loc))
|
| 8 |
+
#loc84 = loc("ks1"(#loc))
|
| 9 |
+
#loc85 = loc("ks2"(#loc))
|
| 10 |
+
#loc86 = loc("ks3"(#loc))
|
| 11 |
+
#loc87 = loc("ks4"(#loc))
|
| 12 |
+
#loc88 = loc("xnumel"(#loc))
|
| 13 |
+
module {
|
| 14 |
+
tt.func public @triton_poi_fused_add_cat_index_mul_neg_slice_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)), %ks4: i64 loc("ks4"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
|
| 15 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc89)
|
| 16 |
+
%xoffset_0 = arith.constant 1024 : i32 loc(#loc90)
|
| 17 |
+
%xoffset_1 = arith.constant 1024 : i32 loc(#loc90)
|
| 18 |
+
%xoffset_2 = arith.muli %xoffset, %xoffset_1 : i32 loc(#loc90)
|
| 19 |
+
%xindex = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc91)
|
| 20 |
+
%xindex_3 = tt.splat %xoffset_2 : i32 -> tensor<1024xi32> loc(#loc92)
|
| 21 |
+
%xindex_4 = arith.addi %xindex_3, %xindex : tensor<1024xi32> loc(#loc92)
|
| 22 |
+
%xmask = tt.splat %xnumel : i32 -> tensor<1024xi32> loc(#loc93)
|
| 23 |
+
%xmask_5 = arith.cmpi slt, %xindex_4, %xmask : tensor<1024xi32> loc(#loc93)
|
| 24 |
+
%x2 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc94)
|
| 25 |
+
%x2_6 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc94)
|
| 26 |
+
%x2_7 = arith.divsi %x2, %x2_6 : tensor<1024xi64> loc(#loc94)
|
| 27 |
+
%x2_8 = tt.splat %ks1 : i64 -> tensor<1024xi64> loc(#loc95)
|
| 28 |
+
%x2_9 = arith.remsi %x2_7, %x2_8 : tensor<1024xi64> loc(#loc95)
|
| 29 |
+
%x0 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc96)
|
| 30 |
+
%x0_10 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc96)
|
| 31 |
+
%x0_11 = arith.remsi %x0, %x0_10 : tensor<1024xi64> loc(#loc96)
|
| 32 |
+
%x5 = arith.extsi %xindex_4 : tensor<1024xi32> to tensor<1024xi64> loc(#loc97)
|
| 33 |
+
%x5_12 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc97)
|
| 34 |
+
%x5_13 = arith.divsi %x5, %x5_12 : tensor<1024xi64> loc(#loc97)
|
| 35 |
+
%tmp0 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc98)
|
| 36 |
+
%tmp0_14 = tt.addptr %tmp0, %xindex_4 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc98)
|
| 37 |
+
%tmp0_15 = tt.load %tmp0_14, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc99)
|
| 38 |
+
%tmp0_16 = arith.extf %tmp0_15 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc100)
|
| 39 |
+
%tmp1 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>> loc(#loc101)
|
| 40 |
+
%tmp1_17 = tt.addptr %tmp1, %x2_9 : tensor<1024x!tt.ptr<i64>>, tensor<1024xi64> loc(#loc101)
|
| 41 |
+
%tmp1_18 = tt.load %tmp1_17, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc102)
|
| 42 |
+
%tmp3 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc103)
|
| 43 |
+
%tmp3_19 = arith.addi %tmp1_18, %tmp3 : tensor<1024xi64> loc(#loc103)
|
| 44 |
+
%tmp4 = arith.constant 0 : i32 loc(#loc104)
|
| 45 |
+
%tmp4_20 = arith.extsi %tmp4 : i32 to i64 loc(#loc104)
|
| 46 |
+
%tmp4_21 = tt.splat %tmp4_20 : i64 -> tensor<1024xi64> loc(#loc104)
|
| 47 |
+
%tmp4_22 = arith.cmpi slt, %tmp1_18, %tmp4_21 : tensor<1024xi64> loc(#loc104)
|
| 48 |
+
%tmp5 = arith.select %tmp4_22, %tmp3_19, %tmp1_18 : tensor<1024xi1>, tensor<1024xi64> loc(#loc105)
|
| 49 |
+
%c0_i32 = arith.constant 0 : i32 loc(#loc18)
|
| 50 |
+
%0 = arith.extsi %c0_i32 : i32 to i64 loc(#loc18)
|
| 51 |
+
%1 = tt.splat %0 : i64 -> tensor<1024xi64> loc(#loc18)
|
| 52 |
+
%2 = arith.cmpi sle, %1, %tmp5 : tensor<1024xi64> loc(#loc18)
|
| 53 |
+
%3 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc19)
|
| 54 |
+
%4 = arith.cmpi slt, %tmp5, %3 : tensor<1024xi64> loc(#loc19)
|
| 55 |
+
%5 = arith.andi %2, %4 : tensor<1024xi1> loc(#loc20)
|
| 56 |
+
%true = arith.constant true loc(#loc21)
|
| 57 |
+
%cst = arith.constant dense<true> : tensor<1024xi1> loc(#loc21)
|
| 58 |
+
%6 = arith.xori %xmask_5, %cst : tensor<1024xi1> loc(#loc21)
|
| 59 |
+
%7 = arith.ori %5, %6 : tensor<1024xi1> loc(#loc22)
|
| 60 |
+
tt.assert %7, "index out of bounds: 0 <= tmp5 < ks2" : tensor<1024xi1> loc(#loc23)
|
| 61 |
+
%tmp7 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc106)
|
| 62 |
+
%tmp7_23 = arith.muli %tmp7, %tmp5 : tensor<1024xi64> loc(#loc106)
|
| 63 |
+
%tmp7_24 = arith.addi %x0_11, %tmp7_23 : tensor<1024xi64> loc(#loc107)
|
| 64 |
+
%tmp7_25 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc108)
|
| 65 |
+
%tmp7_26 = tt.addptr %tmp7_25, %tmp7_24 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc108)
|
| 66 |
+
%tmp7_27 = tt.load %tmp7_26, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc109)
|
| 67 |
+
%tmp7_28 = arith.extf %tmp7_27 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc110)
|
| 68 |
+
%tmp8 = arith.mulf %tmp0_16, %tmp7_28 : tensor<1024xf32> loc(#loc111)
|
| 69 |
+
%tmp10 = arith.constant 0 : i64 loc(#loc112)
|
| 70 |
+
%tmp10_29 = arith.constant dense<0> : tensor<1xi64> loc(#loc112)
|
| 71 |
+
%tmp11 = arith.constant dense<0> : tensor<1024xi64> loc(#loc113)
|
| 72 |
+
%tmp11_30 = arith.cmpi sge, %x0_11, %tmp11 : tensor<1024xi64> loc(#loc113)
|
| 73 |
+
%tmp12 = arith.constant 2 : i32 loc(#loc114)
|
| 74 |
+
%tmp12_31 = arith.constant 2 : i64 loc(#loc114)
|
| 75 |
+
%tmp12_32 = arith.divsi %ks3, %tmp12_31 : i64 loc(#loc114)
|
| 76 |
+
%tmp12_33 = arith.constant -1 : i32 loc(#loc115)
|
| 77 |
+
%tmp12_34 = arith.constant -1 : i64 loc(#loc115)
|
| 78 |
+
%tmp12_35 = arith.muli %tmp12_34, %tmp12_32 : i64 loc(#loc115)
|
| 79 |
+
%tmp12_36 = arith.addi %ks3, %tmp12_35 : i64 loc(#loc116)
|
| 80 |
+
%tmp13 = tt.splat %tmp12_36 : i64 -> tensor<1024xi64> loc(#loc117)
|
| 81 |
+
%tmp13_37 = arith.cmpi slt, %x0_11, %tmp13 : tensor<1024xi64> loc(#loc117)
|
| 82 |
+
%tmp14 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc118)
|
| 83 |
+
%tmp14_38 = arith.muli %tmp14, %x5_13 : tensor<1024xi64> loc(#loc118)
|
| 84 |
+
%tmp14_39 = arith.constant 2 : i32 loc(#loc119)
|
| 85 |
+
%tmp14_40 = arith.constant 2 : i64 loc(#loc119)
|
| 86 |
+
%tmp14_41 = arith.divsi %ks3, %tmp14_40 : i64 loc(#loc119)
|
| 87 |
+
%tmp14_42 = tt.splat %tmp14_41 : i64 -> tensor<1024xi64> loc(#loc120)
|
| 88 |
+
%tmp14_43 = arith.addi %tmp14_38, %tmp14_42 : tensor<1024xi64> loc(#loc120)
|
| 89 |
+
%tmp14_44 = arith.addi %tmp14_43, %x0_11 : tensor<1024xi64> loc(#loc121)
|
| 90 |
+
%tmp14_45 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc122)
|
| 91 |
+
%tmp14_46 = tt.addptr %tmp14_45, %tmp14_44 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc122)
|
| 92 |
+
%tmp14_47 = arith.andi %tmp13_37, %xmask_5 : tensor<1024xi1> loc(#loc123)
|
| 93 |
+
%tmp14_48 = arith.constant 0.000000e+00 : f32 loc(#loc124)
|
| 94 |
+
%tmp14_49 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc124)
|
| 95 |
+
%tmp14_50 = arith.truncf %tmp14_49 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc124)
|
| 96 |
+
%tmp14_51 = tt.load %tmp14_46, %tmp14_47, %tmp14_50 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc124)
|
| 97 |
+
%tmp14_52 = arith.extf %tmp14_51 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc125)
|
| 98 |
+
%tmp15 = arith.constant 0.000000e+00 : f32 loc(#loc126)
|
| 99 |
+
%tmp15_53 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc126)
|
| 100 |
+
%tmp15_54 = arith.subf %tmp15_53, %tmp14_52 : tensor<1024xf32> loc(#loc126)
|
| 101 |
+
%tmp16 = arith.constant 0.000000e+00 : f32 loc(#loc127)
|
| 102 |
+
%tmp16_55 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc127)
|
| 103 |
+
%tmp17 = arith.select %tmp13_37, %tmp15_54, %tmp16_55 : tensor<1024xi1>, tensor<1024xf32> loc(#loc128)
|
| 104 |
+
%tmp18 = tt.splat %tmp12_36 : i64 -> tensor<1024xi64> loc(#loc129)
|
| 105 |
+
%tmp18_56 = arith.cmpi sge, %x0_11, %tmp18 : tensor<1024xi64> loc(#loc129)
|
| 106 |
+
%tmp20 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc130)
|
| 107 |
+
%tmp20_57 = arith.cmpi slt, %x0_11, %tmp20 : tensor<1024xi64> loc(#loc130)
|
| 108 |
+
%tmp21 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc131)
|
| 109 |
+
%tmp21_58 = arith.muli %tmp21, %x5_13 : tensor<1024xi64> loc(#loc131)
|
| 110 |
+
%tmp21_59 = arith.constant -1 : i32 loc(#loc132)
|
| 111 |
+
%tmp21_60 = arith.constant -1 : i64 loc(#loc132)
|
| 112 |
+
%tmp21_61 = arith.muli %tmp21_60, %ks3 : i64 loc(#loc132)
|
| 113 |
+
%tmp21_62 = tt.splat %tmp21_61 : i64 -> tensor<1024xi64> loc(#loc133)
|
| 114 |
+
%tmp21_63 = arith.addi %x0_11, %tmp21_62 : tensor<1024xi64> loc(#loc133)
|
| 115 |
+
%tmp21_64 = arith.constant 2 : i32 loc(#loc134)
|
| 116 |
+
%tmp21_65 = arith.constant 2 : i64 loc(#loc134)
|
| 117 |
+
%tmp21_66 = arith.divsi %ks3, %tmp21_65 : i64 loc(#loc134)
|
| 118 |
+
%tmp21_67 = tt.splat %tmp21_66 : i64 -> tensor<1024xi64> loc(#loc135)
|
| 119 |
+
%tmp21_68 = arith.addi %tmp21_63, %tmp21_67 : tensor<1024xi64> loc(#loc135)
|
| 120 |
+
%tmp21_69 = arith.addi %tmp21_58, %tmp21_68 : tensor<1024xi64> loc(#loc136)
|
| 121 |
+
%tmp21_70 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc137)
|
| 122 |
+
%tmp21_71 = tt.addptr %tmp21_70, %tmp21_69 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc137)
|
| 123 |
+
%tmp21_72 = arith.andi %tmp18_56, %xmask_5 : tensor<1024xi1> loc(#loc138)
|
| 124 |
+
%tmp21_73 = arith.constant 0.000000e+00 : f32 loc(#loc139)
|
| 125 |
+
%tmp21_74 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc139)
|
| 126 |
+
%tmp21_75 = arith.truncf %tmp21_74 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc139)
|
| 127 |
+
%tmp21_76 = tt.load %tmp21_71, %tmp21_72, %tmp21_75 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc139)
|
| 128 |
+
%tmp21_77 = arith.extf %tmp21_76 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc140)
|
| 129 |
+
%tmp22 = arith.select %tmp13_37, %tmp17, %tmp21_77 : tensor<1024xi1>, tensor<1024xf32> loc(#loc141)
|
| 130 |
+
%tmp24 = tt.splat %ks4 : i64 -> tensor<1024xi64> loc(#loc142)
|
| 131 |
+
%tmp24_78 = arith.addi %tmp1_18, %tmp24 : tensor<1024xi64> loc(#loc142)
|
| 132 |
+
%tmp25 = arith.select %tmp4_22, %tmp24_78, %tmp1_18 : tensor<1024xi1>, tensor<1024xi64> loc(#loc143)
|
| 133 |
+
%c0_i32_79 = arith.constant 0 : i32 loc(#loc62)
|
| 134 |
+
%8 = arith.extsi %c0_i32_79 : i32 to i64 loc(#loc62)
|
| 135 |
+
%9 = tt.splat %8 : i64 -> tensor<1024xi64> loc(#loc62)
|
| 136 |
+
%10 = arith.cmpi sle, %9, %tmp25 : tensor<1024xi64> loc(#loc62)
|
| 137 |
+
%11 = tt.splat %ks4 : i64 -> tensor<1024xi64> loc(#loc63)
|
| 138 |
+
%12 = arith.cmpi slt, %tmp25, %11 : tensor<1024xi64> loc(#loc63)
|
| 139 |
+
%13 = arith.andi %10, %12 : tensor<1024xi1> loc(#loc64)
|
| 140 |
+
%true_80 = arith.constant true loc(#loc65)
|
| 141 |
+
%cst_81 = arith.constant dense<true> : tensor<1024xi1> loc(#loc65)
|
| 142 |
+
%14 = arith.xori %xmask_5, %cst_81 : tensor<1024xi1> loc(#loc65)
|
| 143 |
+
%15 = arith.ori %13, %14 : tensor<1024xi1> loc(#loc66)
|
| 144 |
+
tt.assert %15, "index out of bounds: 0 <= tmp25 < ks4" : tensor<1024xi1> loc(#loc67)
|
| 145 |
+
%tmp27 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc144)
|
| 146 |
+
%tmp27_82 = arith.muli %tmp27, %tmp25 : tensor<1024xi64> loc(#loc144)
|
| 147 |
+
%tmp27_83 = arith.addi %x0_11, %tmp27_82 : tensor<1024xi64> loc(#loc145)
|
| 148 |
+
%tmp27_84 = tt.splat %in_ptr3 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc146)
|
| 149 |
+
%tmp27_85 = tt.addptr %tmp27_84, %tmp27_83 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc146)
|
| 150 |
+
%tmp27_86 = tt.load %tmp27_85, %xmask_5 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc147)
|
| 151 |
+
%tmp27_87 = arith.extf %tmp27_86 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc148)
|
| 152 |
+
%tmp28 = arith.mulf %tmp22, %tmp27_87 : tensor<1024xf32> loc(#loc149)
|
| 153 |
+
%tmp29 = arith.addf %tmp8, %tmp28 : tensor<1024xf32> loc(#loc150)
|
| 154 |
+
%16 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc75)
|
| 155 |
+
%17 = tt.addptr %16, %xindex_4 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc75)
|
| 156 |
+
%18 = arith.truncf %tmp29 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc76)
|
| 157 |
+
tt.store %17, %18, %xmask_5 : tensor<1024x!tt.ptr<bf16>> loc(#loc76)
|
| 158 |
+
tt.return loc(#loc77)
|
| 159 |
+
} loc(#loc)
|
| 160 |
+
} loc(#loc)
|
| 161 |
+
#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":19:28)
|
| 162 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":19:33)
|
| 163 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":20:36)
|
| 164 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":20:23)
|
| 165 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":21:21)
|
| 166 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":23:21)
|
| 167 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":23:28)
|
| 168 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":24:19)
|
| 169 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":25:19)
|
| 170 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:30)
|
| 171 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:35)
|
| 172 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:75)
|
| 173 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":27:30)
|
| 174 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":27:35)
|
| 175 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":29:18)
|
| 176 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":30:18)
|
| 177 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":31:32)
|
| 178 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:28)
|
| 179 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:44)
|
| 180 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:37)
|
| 181 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:54)
|
| 182 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:52)
|
| 183 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:62)
|
| 184 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:39)
|
| 185 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:35)
|
| 186 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:30)
|
| 187 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:46)
|
| 188 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:86)
|
| 189 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":34:18)
|
| 190 |
+
#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":36:28)
|
| 191 |
+
#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":37:20)
|
| 192 |
+
#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:31)
|
| 193 |
+
#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:24)
|
| 194 |
+
#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:18)
|
| 195 |
+
#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":39:19)
|
| 196 |
+
#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:35)
|
| 197 |
+
#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:48)
|
| 198 |
+
#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:41)
|
| 199 |
+
#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:54)
|
| 200 |
+
#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:31)
|
| 201 |
+
#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:68)
|
| 202 |
+
#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:60)
|
| 203 |
+
#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:119)
|
| 204 |
+
#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":41:13)
|
| 205 |
+
#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":42:38)
|
| 206 |
+
#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":43:35)
|
| 207 |
+
#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":44:20)
|
| 208 |
+
#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":46:19)
|
| 209 |
+
#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:35)
|
| 210 |
+
#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:52)
|
| 211 |
+
#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:47)
|
| 212 |
+
#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:67)
|
| 213 |
+
#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:60)
|
| 214 |
+
#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:41)
|
| 215 |
+
#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:31)
|
| 216 |
+
#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:81)
|
| 217 |
+
#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:73)
|
| 218 |
+
#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:132)
|
| 219 |
+
#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":48:35)
|
| 220 |
+
#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":50:19)
|
| 221 |
+
#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":51:34)
|
| 222 |
+
#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:28)
|
| 223 |
+
#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:46)
|
| 224 |
+
#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:38)
|
| 225 |
+
#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:56)
|
| 226 |
+
#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:54)
|
| 227 |
+
#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:64)
|
| 228 |
+
#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:40)
|
| 229 |
+
#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:36)
|
| 230 |
+
#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:31)
|
| 231 |
+
#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:48)
|
| 232 |
+
#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:88)
|
| 233 |
+
#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":54:20)
|
| 234 |
+
#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":55:19)
|
| 235 |
+
#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:25)
|
| 236 |
+
#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:37)
|
| 237 |
+
#loc77 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:4)
|
| 238 |
+
#loc89 = loc("xoffset"(#loc1))
|
| 239 |
+
#loc90 = loc("xoffset"(#loc2))
|
| 240 |
+
#loc91 = loc("xindex"(#loc3))
|
| 241 |
+
#loc92 = loc("xindex"(#loc4))
|
| 242 |
+
#loc93 = loc("xmask"(#loc5))
|
| 243 |
+
#loc94 = loc("x2"(#loc6))
|
| 244 |
+
#loc95 = loc("x2"(#loc7))
|
| 245 |
+
#loc96 = loc("x0"(#loc8))
|
| 246 |
+
#loc97 = loc("x5"(#loc9))
|
| 247 |
+
#loc98 = loc("tmp0"(#loc10))
|
| 248 |
+
#loc99 = loc("tmp0"(#loc11))
|
| 249 |
+
#loc100 = loc("tmp0"(#loc12))
|
| 250 |
+
#loc101 = loc("tmp1"(#loc13))
|
| 251 |
+
#loc102 = loc("tmp1"(#loc14))
|
| 252 |
+
#loc103 = loc("tmp3"(#loc15))
|
| 253 |
+
#loc104 = loc("tmp4"(#loc16))
|
| 254 |
+
#loc105 = loc("tmp5"(#loc17))
|
| 255 |
+
#loc106 = loc("tmp7"(#loc24))
|
| 256 |
+
#loc107 = loc("tmp7"(#loc25))
|
| 257 |
+
#loc108 = loc("tmp7"(#loc26))
|
| 258 |
+
#loc109 = loc("tmp7"(#loc27))
|
| 259 |
+
#loc110 = loc("tmp7"(#loc28))
|
| 260 |
+
#loc111 = loc("tmp8"(#loc29))
|
| 261 |
+
#loc112 = loc("tmp10"(#loc30))
|
| 262 |
+
#loc113 = loc("tmp11"(#loc31))
|
| 263 |
+
#loc114 = loc("tmp12"(#loc32))
|
| 264 |
+
#loc115 = loc("tmp12"(#loc33))
|
| 265 |
+
#loc116 = loc("tmp12"(#loc34))
|
| 266 |
+
#loc117 = loc("tmp13"(#loc35))
|
| 267 |
+
#loc118 = loc("tmp14"(#loc36))
|
| 268 |
+
#loc119 = loc("tmp14"(#loc37))
|
| 269 |
+
#loc120 = loc("tmp14"(#loc38))
|
| 270 |
+
#loc121 = loc("tmp14"(#loc39))
|
| 271 |
+
#loc122 = loc("tmp14"(#loc40))
|
| 272 |
+
#loc123 = loc("tmp14"(#loc41))
|
| 273 |
+
#loc124 = loc("tmp14"(#loc42))
|
| 274 |
+
#loc125 = loc("tmp14"(#loc43))
|
| 275 |
+
#loc126 = loc("tmp15"(#loc44))
|
| 276 |
+
#loc127 = loc("tmp16"(#loc45))
|
| 277 |
+
#loc128 = loc("tmp17"(#loc46))
|
| 278 |
+
#loc129 = loc("tmp18"(#loc47))
|
| 279 |
+
#loc130 = loc("tmp20"(#loc48))
|
| 280 |
+
#loc131 = loc("tmp21"(#loc49))
|
| 281 |
+
#loc132 = loc("tmp21"(#loc50))
|
| 282 |
+
#loc133 = loc("tmp21"(#loc51))
|
| 283 |
+
#loc134 = loc("tmp21"(#loc52))
|
| 284 |
+
#loc135 = loc("tmp21"(#loc53))
|
| 285 |
+
#loc136 = loc("tmp21"(#loc54))
|
| 286 |
+
#loc137 = loc("tmp21"(#loc55))
|
| 287 |
+
#loc138 = loc("tmp21"(#loc56))
|
| 288 |
+
#loc139 = loc("tmp21"(#loc57))
|
| 289 |
+
#loc140 = loc("tmp21"(#loc58))
|
| 290 |
+
#loc141 = loc("tmp22"(#loc59))
|
| 291 |
+
#loc142 = loc("tmp24"(#loc60))
|
| 292 |
+
#loc143 = loc("tmp25"(#loc61))
|
| 293 |
+
#loc144 = loc("tmp27"(#loc68))
|
| 294 |
+
#loc145 = loc("tmp27"(#loc69))
|
| 295 |
+
#loc146 = loc("tmp27"(#loc70))
|
| 296 |
+
#loc147 = loc("tmp27"(#loc71))
|
| 297 |
+
#loc148 = loc("tmp27"(#loc72))
|
| 298 |
+
#loc149 = loc("tmp28"(#loc73))
|
| 299 |
+
#loc150 = loc("tmp29"(#loc74))
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttgir
ADDED
|
@@ -0,0 +1,232 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#blocked = #ttg.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
|
| 2 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":18:0)
|
| 3 |
+
#loc70 = loc("in_ptr0"(#loc))
|
| 4 |
+
#loc71 = loc("in_ptr1"(#loc))
|
| 5 |
+
#loc72 = loc("in_ptr2"(#loc))
|
| 6 |
+
#loc73 = loc("in_ptr3"(#loc))
|
| 7 |
+
#loc74 = loc("out_ptr0"(#loc))
|
| 8 |
+
#loc75 = loc("ks0"(#loc))
|
| 9 |
+
#loc76 = loc("ks1"(#loc))
|
| 10 |
+
#loc77 = loc("ks2"(#loc))
|
| 11 |
+
#loc78 = loc("ks3"(#loc))
|
| 12 |
+
#loc79 = loc("ks4"(#loc))
|
| 13 |
+
#loc80 = loc("xnumel"(#loc))
|
| 14 |
+
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
|
| 15 |
+
tt.func public @triton_poi_fused_add_cat_index_mul_neg_slice_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)), %ks4: i64 loc("ks4"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
|
| 16 |
+
%cst = arith.constant dense<true> : tensor<1024xi1, #blocked> loc(#loc1)
|
| 17 |
+
%c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
|
| 18 |
+
%cst_0 = arith.constant dense<0.000000e+00> : tensor<1024xbf16, #blocked> loc(#loc1)
|
| 19 |
+
%c2_i64 = arith.constant 2 : i64 loc(#loc1)
|
| 20 |
+
%c-1_i64 = arith.constant -1 : i64 loc(#loc1)
|
| 21 |
+
%cst_1 = arith.constant dense<0> : tensor<1024xi64, #blocked> loc(#loc1)
|
| 22 |
+
%cst_2 = arith.constant dense<0.000000e+00> : tensor<1024xf32, #blocked> loc(#loc1)
|
| 23 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc81)
|
| 24 |
+
%xoffset_3 = arith.muli %xoffset, %c1024_i32 : i32 loc(#loc82)
|
| 25 |
+
%xindex = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #blocked> loc(#loc83)
|
| 26 |
+
%xindex_4 = tt.splat %xoffset_3 : i32 -> tensor<1024xi32, #blocked> loc(#loc84)
|
| 27 |
+
%xindex_5 = arith.addi %xindex_4, %xindex : tensor<1024xi32, #blocked> loc(#loc84)
|
| 28 |
+
%xmask = tt.splat %xnumel : i32 -> tensor<1024xi32, #blocked> loc(#loc85)
|
| 29 |
+
%xmask_6 = arith.cmpi slt, %xindex_5, %xmask : tensor<1024xi32, #blocked> loc(#loc85)
|
| 30 |
+
%x2 = arith.extsi %xindex_5 : tensor<1024xi32, #blocked> to tensor<1024xi64, #blocked> loc(#loc86)
|
| 31 |
+
%x2_7 = tt.splat %ks0 : i64 -> tensor<1024xi64, #blocked> loc(#loc86)
|
| 32 |
+
%x2_8 = arith.divsi %x2, %x2_7 : tensor<1024xi64, #blocked> loc(#loc86)
|
| 33 |
+
%x2_9 = tt.splat %ks1 : i64 -> tensor<1024xi64, #blocked> loc(#loc87)
|
| 34 |
+
%x2_10 = arith.remsi %x2_8, %x2_9 : tensor<1024xi64, #blocked> loc(#loc87)
|
| 35 |
+
%x0 = tt.splat %ks3 : i64 -> tensor<1024xi64, #blocked> loc(#loc88)
|
| 36 |
+
%x0_11 = arith.remsi %x2, %x0 : tensor<1024xi64, #blocked> loc(#loc88)
|
| 37 |
+
%x5 = arith.divsi %x2, %x0 : tensor<1024xi64, #blocked> loc(#loc89)
|
| 38 |
+
%tmp0 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc90)
|
| 39 |
+
%tmp0_12 = tt.addptr %tmp0, %xindex_5 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi32, #blocked> loc(#loc90)
|
| 40 |
+
%tmp0_13 = tt.load %tmp0_12, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc91)
|
| 41 |
+
%tmp0_14 = arith.extf %tmp0_13 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc92)
|
| 42 |
+
%tmp1 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>, #blocked> loc(#loc93)
|
| 43 |
+
%tmp1_15 = tt.addptr %tmp1, %x2_10 : tensor<1024x!tt.ptr<i64>, #blocked>, tensor<1024xi64, #blocked> loc(#loc93)
|
| 44 |
+
%tmp1_16 = tt.load %tmp1_15, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>, #blocked> loc(#loc94)
|
| 45 |
+
%tmp3 = tt.splat %ks2 : i64 -> tensor<1024xi64, #blocked> loc(#loc95)
|
| 46 |
+
%tmp3_17 = arith.addi %tmp1_16, %tmp3 : tensor<1024xi64, #blocked> loc(#loc95)
|
| 47 |
+
%tmp4 = arith.cmpi slt, %tmp1_16, %cst_1 : tensor<1024xi64, #blocked> loc(#loc96)
|
| 48 |
+
%tmp5 = arith.select %tmp4, %tmp3_17, %tmp1_16 : tensor<1024xi1, #blocked>, tensor<1024xi64, #blocked> loc(#loc97)
|
| 49 |
+
%0 = arith.cmpi sge, %tmp5, %cst_1 : tensor<1024xi64, #blocked> loc(#loc19)
|
| 50 |
+
%1 = arith.cmpi slt, %tmp5, %tmp3 : tensor<1024xi64, #blocked> loc(#loc20)
|
| 51 |
+
%2 = arith.andi %0, %1 : tensor<1024xi1, #blocked> loc(#loc21)
|
| 52 |
+
%3 = arith.xori %xmask_6, %cst : tensor<1024xi1, #blocked> loc(#loc22)
|
| 53 |
+
%4 = arith.ori %2, %3 : tensor<1024xi1, #blocked> loc(#loc23)
|
| 54 |
+
tt.assert %4, "index out of bounds: 0 <= tmp5 < ks2" : tensor<1024xi1, #blocked> loc(#loc24)
|
| 55 |
+
%tmp7 = arith.muli %x0, %tmp5 : tensor<1024xi64, #blocked> loc(#loc98)
|
| 56 |
+
%tmp7_18 = arith.addi %x0_11, %tmp7 : tensor<1024xi64, #blocked> loc(#loc99)
|
| 57 |
+
%tmp7_19 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc100)
|
| 58 |
+
%tmp7_20 = tt.addptr %tmp7_19, %tmp7_18 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc100)
|
| 59 |
+
%tmp7_21 = tt.load %tmp7_20, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc101)
|
| 60 |
+
%tmp7_22 = arith.extf %tmp7_21 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc102)
|
| 61 |
+
%tmp8 = arith.mulf %tmp0_14, %tmp7_22 : tensor<1024xf32, #blocked> loc(#loc103)
|
| 62 |
+
%tmp12 = arith.divsi %ks3, %c2_i64 : i64 loc(#loc104)
|
| 63 |
+
%tmp12_23 = arith.subi %ks3, %tmp12 : i64 loc(#loc105)
|
| 64 |
+
%tmp13 = tt.splat %tmp12_23 : i64 -> tensor<1024xi64, #blocked> loc(#loc106)
|
| 65 |
+
%tmp13_24 = arith.cmpi slt, %x0_11, %tmp13 : tensor<1024xi64, #blocked> loc(#loc106)
|
| 66 |
+
%tmp14 = arith.muli %x0, %x5 : tensor<1024xi64, #blocked> loc(#loc107)
|
| 67 |
+
%tmp14_25 = tt.splat %tmp12 : i64 -> tensor<1024xi64, #blocked> loc(#loc108)
|
| 68 |
+
%tmp14_26 = arith.addi %tmp14, %tmp14_25 : tensor<1024xi64, #blocked> loc(#loc108)
|
| 69 |
+
%tmp14_27 = arith.addi %tmp14_26, %x0_11 : tensor<1024xi64, #blocked> loc(#loc109)
|
| 70 |
+
%tmp14_28 = tt.addptr %tmp0, %tmp14_27 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc110)
|
| 71 |
+
%tmp14_29 = arith.andi %tmp13_24, %xmask_6 : tensor<1024xi1, #blocked> loc(#loc111)
|
| 72 |
+
%tmp14_30 = tt.load %tmp14_28, %tmp14_29, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc112)
|
| 73 |
+
%tmp14_31 = arith.extf %tmp14_30 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc113)
|
| 74 |
+
%tmp15 = arith.subf %cst_2, %tmp14_31 : tensor<1024xf32, #blocked> loc(#loc114)
|
| 75 |
+
%tmp18 = arith.cmpi sge, %x0_11, %tmp13 : tensor<1024xi64, #blocked> loc(#loc115)
|
| 76 |
+
%tmp21 = arith.muli %ks3, %c-1_i64 : i64 loc(#loc116)
|
| 77 |
+
%tmp21_32 = tt.splat %tmp21 : i64 -> tensor<1024xi64, #blocked> loc(#loc117)
|
| 78 |
+
%tmp21_33 = arith.addi %x0_11, %tmp21_32 : tensor<1024xi64, #blocked> loc(#loc117)
|
| 79 |
+
%tmp21_34 = arith.addi %tmp21_33, %tmp14_25 : tensor<1024xi64, #blocked> loc(#loc118)
|
| 80 |
+
%tmp21_35 = arith.addi %tmp14, %tmp21_34 : tensor<1024xi64, #blocked> loc(#loc119)
|
| 81 |
+
%tmp21_36 = tt.addptr %tmp0, %tmp21_35 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc120)
|
| 82 |
+
%tmp21_37 = arith.andi %tmp18, %xmask_6 : tensor<1024xi1, #blocked> loc(#loc121)
|
| 83 |
+
%tmp21_38 = tt.load %tmp21_36, %tmp21_37, %cst_0 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc122)
|
| 84 |
+
%tmp21_39 = arith.extf %tmp21_38 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc123)
|
| 85 |
+
%tmp22 = arith.select %tmp13_24, %tmp15, %tmp21_39 : tensor<1024xi1, #blocked>, tensor<1024xf32, #blocked> loc(#loc135)
|
| 86 |
+
%tmp24 = tt.splat %ks4 : i64 -> tensor<1024xi64, #blocked> loc(#loc126)
|
| 87 |
+
%tmp24_40 = arith.addi %tmp1_16, %tmp24 : tensor<1024xi64, #blocked> loc(#loc126)
|
| 88 |
+
%tmp25 = arith.select %tmp4, %tmp24_40, %tmp1_16 : tensor<1024xi1, #blocked>, tensor<1024xi64, #blocked> loc(#loc127)
|
| 89 |
+
%5 = arith.cmpi sge, %tmp25, %cst_1 : tensor<1024xi64, #blocked> loc(#loc55)
|
| 90 |
+
%6 = arith.cmpi slt, %tmp25, %tmp24 : tensor<1024xi64, #blocked> loc(#loc56)
|
| 91 |
+
%7 = arith.andi %5, %6 : tensor<1024xi1, #blocked> loc(#loc57)
|
| 92 |
+
%8 = arith.ori %7, %3 : tensor<1024xi1, #blocked> loc(#loc58)
|
| 93 |
+
tt.assert %8, "index out of bounds: 0 <= tmp25 < ks4" : tensor<1024xi1, #blocked> loc(#loc59)
|
| 94 |
+
%tmp27 = arith.muli %x0, %tmp25 : tensor<1024xi64, #blocked> loc(#loc128)
|
| 95 |
+
%tmp27_41 = arith.addi %x0_11, %tmp27 : tensor<1024xi64, #blocked> loc(#loc129)
|
| 96 |
+
%tmp27_42 = tt.splat %in_ptr3 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc130)
|
| 97 |
+
%tmp27_43 = tt.addptr %tmp27_42, %tmp27_41 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi64, #blocked> loc(#loc130)
|
| 98 |
+
%tmp27_44 = tt.load %tmp27_43, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc131)
|
| 99 |
+
%tmp27_45 = arith.extf %tmp27_44 : tensor<1024xbf16, #blocked> to tensor<1024xf32, #blocked> loc(#loc132)
|
| 100 |
+
%tmp28 = arith.mulf %tmp22, %tmp27_45 : tensor<1024xf32, #blocked> loc(#loc133)
|
| 101 |
+
%tmp29 = arith.addf %tmp8, %tmp28 : tensor<1024xf32, #blocked> loc(#loc134)
|
| 102 |
+
%9 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc67)
|
| 103 |
+
%10 = tt.addptr %9, %xindex_5 : tensor<1024x!tt.ptr<bf16>, #blocked>, tensor<1024xi32, #blocked> loc(#loc67)
|
| 104 |
+
%11 = arith.truncf %tmp29 : tensor<1024xf32, #blocked> to tensor<1024xbf16, #blocked> loc(#loc68)
|
| 105 |
+
tt.store %10, %11, %xmask_6 : tensor<1024x!tt.ptr<bf16>, #blocked> loc(#loc68)
|
| 106 |
+
tt.return loc(#loc69)
|
| 107 |
+
} loc(#loc)
|
| 108 |
+
} loc(#loc)
|
| 109 |
+
#loc1 = loc(unknown)
|
| 110 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":19:28)
|
| 111 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":19:33)
|
| 112 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":20:36)
|
| 113 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":20:23)
|
| 114 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":21:21)
|
| 115 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":23:21)
|
| 116 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":23:28)
|
| 117 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":24:19)
|
| 118 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":25:19)
|
| 119 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:30)
|
| 120 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:35)
|
| 121 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:75)
|
| 122 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":27:30)
|
| 123 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":27:35)
|
| 124 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":29:18)
|
| 125 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":30:18)
|
| 126 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":31:32)
|
| 127 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:28)
|
| 128 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:44)
|
| 129 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:37)
|
| 130 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:54)
|
| 131 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:52)
|
| 132 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:62)
|
| 133 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:39)
|
| 134 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:35)
|
| 135 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:30)
|
| 136 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:46)
|
| 137 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:86)
|
| 138 |
+
#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":34:18)
|
| 139 |
+
#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:31)
|
| 140 |
+
#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:18)
|
| 141 |
+
#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":39:19)
|
| 142 |
+
#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:35)
|
| 143 |
+
#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:41)
|
| 144 |
+
#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:54)
|
| 145 |
+
#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:31)
|
| 146 |
+
#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:68)
|
| 147 |
+
#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:60)
|
| 148 |
+
#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:119)
|
| 149 |
+
#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":41:13)
|
| 150 |
+
#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":44:20)
|
| 151 |
+
#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:52)
|
| 152 |
+
#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:47)
|
| 153 |
+
#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:60)
|
| 154 |
+
#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:41)
|
| 155 |
+
#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:31)
|
| 156 |
+
#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:81)
|
| 157 |
+
#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:73)
|
| 158 |
+
#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:132)
|
| 159 |
+
#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":48:35)
|
| 160 |
+
#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":43:35)
|
| 161 |
+
#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":50:19)
|
| 162 |
+
#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":51:34)
|
| 163 |
+
#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:28)
|
| 164 |
+
#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:46)
|
| 165 |
+
#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:38)
|
| 166 |
+
#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:54)
|
| 167 |
+
#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:64)
|
| 168 |
+
#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:40)
|
| 169 |
+
#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:36)
|
| 170 |
+
#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:31)
|
| 171 |
+
#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:48)
|
| 172 |
+
#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:88)
|
| 173 |
+
#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":54:20)
|
| 174 |
+
#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":55:19)
|
| 175 |
+
#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:25)
|
| 176 |
+
#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:37)
|
| 177 |
+
#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:4)
|
| 178 |
+
#loc81 = loc("xoffset"(#loc2))
|
| 179 |
+
#loc82 = loc("xoffset"(#loc3))
|
| 180 |
+
#loc83 = loc("xindex"(#loc4))
|
| 181 |
+
#loc84 = loc("xindex"(#loc5))
|
| 182 |
+
#loc85 = loc("xmask"(#loc6))
|
| 183 |
+
#loc86 = loc("x2"(#loc7))
|
| 184 |
+
#loc87 = loc("x2"(#loc8))
|
| 185 |
+
#loc88 = loc("x0"(#loc9))
|
| 186 |
+
#loc89 = loc("x5"(#loc10))
|
| 187 |
+
#loc90 = loc("tmp0"(#loc11))
|
| 188 |
+
#loc91 = loc("tmp0"(#loc12))
|
| 189 |
+
#loc92 = loc("tmp0"(#loc13))
|
| 190 |
+
#loc93 = loc("tmp1"(#loc14))
|
| 191 |
+
#loc94 = loc("tmp1"(#loc15))
|
| 192 |
+
#loc95 = loc("tmp3"(#loc16))
|
| 193 |
+
#loc96 = loc("tmp4"(#loc17))
|
| 194 |
+
#loc97 = loc("tmp5"(#loc18))
|
| 195 |
+
#loc98 = loc("tmp7"(#loc25))
|
| 196 |
+
#loc99 = loc("tmp7"(#loc26))
|
| 197 |
+
#loc100 = loc("tmp7"(#loc27))
|
| 198 |
+
#loc101 = loc("tmp7"(#loc28))
|
| 199 |
+
#loc102 = loc("tmp7"(#loc29))
|
| 200 |
+
#loc103 = loc("tmp8"(#loc30))
|
| 201 |
+
#loc104 = loc("tmp12"(#loc31))
|
| 202 |
+
#loc105 = loc("tmp12"(#loc32))
|
| 203 |
+
#loc106 = loc("tmp13"(#loc33))
|
| 204 |
+
#loc107 = loc("tmp14"(#loc34))
|
| 205 |
+
#loc108 = loc("tmp14"(#loc35))
|
| 206 |
+
#loc109 = loc("tmp14"(#loc36))
|
| 207 |
+
#loc110 = loc("tmp14"(#loc37))
|
| 208 |
+
#loc111 = loc("tmp14"(#loc38))
|
| 209 |
+
#loc112 = loc("tmp14"(#loc39))
|
| 210 |
+
#loc113 = loc("tmp14"(#loc40))
|
| 211 |
+
#loc114 = loc("tmp15"(#loc41))
|
| 212 |
+
#loc115 = loc("tmp18"(#loc42))
|
| 213 |
+
#loc116 = loc("tmp21"(#loc43))
|
| 214 |
+
#loc117 = loc("tmp21"(#loc44))
|
| 215 |
+
#loc118 = loc("tmp21"(#loc45))
|
| 216 |
+
#loc119 = loc("tmp21"(#loc46))
|
| 217 |
+
#loc120 = loc("tmp21"(#loc47))
|
| 218 |
+
#loc121 = loc("tmp21"(#loc48))
|
| 219 |
+
#loc122 = loc("tmp21"(#loc49))
|
| 220 |
+
#loc123 = loc("tmp21"(#loc50))
|
| 221 |
+
#loc124 = loc("tmp22"(#loc51))
|
| 222 |
+
#loc125 = loc("tmp17"(#loc52))
|
| 223 |
+
#loc126 = loc("tmp24"(#loc53))
|
| 224 |
+
#loc127 = loc("tmp25"(#loc54))
|
| 225 |
+
#loc128 = loc("tmp27"(#loc60))
|
| 226 |
+
#loc129 = loc("tmp27"(#loc61))
|
| 227 |
+
#loc130 = loc("tmp27"(#loc62))
|
| 228 |
+
#loc131 = loc("tmp27"(#loc63))
|
| 229 |
+
#loc132 = loc("tmp27"(#loc64))
|
| 230 |
+
#loc133 = loc("tmp28"(#loc65))
|
| 231 |
+
#loc134 = loc("tmp29"(#loc66))
|
| 232 |
+
#loc135 = loc(fused[#loc124, #loc125])
|
SpecForge-ext/cache/compiled_kernels/triton/1/2TU6ZCF6AOXLWQQED5J7FS5ZXMYK7TIOQ6T2MLB767275BROXJCA/triton_poi_fused_add_cat_index_mul_neg_slice_squeeze_unsqueeze_0.ttir
ADDED
|
@@ -0,0 +1,231 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":18:0)
|
| 2 |
+
#loc70 = loc("in_ptr0"(#loc))
|
| 3 |
+
#loc71 = loc("in_ptr1"(#loc))
|
| 4 |
+
#loc72 = loc("in_ptr2"(#loc))
|
| 5 |
+
#loc73 = loc("in_ptr3"(#loc))
|
| 6 |
+
#loc74 = loc("out_ptr0"(#loc))
|
| 7 |
+
#loc75 = loc("ks0"(#loc))
|
| 8 |
+
#loc76 = loc("ks1"(#loc))
|
| 9 |
+
#loc77 = loc("ks2"(#loc))
|
| 10 |
+
#loc78 = loc("ks3"(#loc))
|
| 11 |
+
#loc79 = loc("ks4"(#loc))
|
| 12 |
+
#loc80 = loc("xnumel"(#loc))
|
| 13 |
+
module {
|
| 14 |
+
tt.func public @triton_poi_fused_add_cat_index_mul_neg_slice_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)), %ks4: i64 loc("ks4"(#loc)), %xnumel: i32 loc("xnumel"(#loc))) attributes {noinline = false} {
|
| 15 |
+
%cst = arith.constant dense<0.000000e+00> : tensor<1024xbf16> loc(#loc1)
|
| 16 |
+
%cst_0 = arith.constant dense<0> : tensor<1024xi64> loc(#loc1)
|
| 17 |
+
%cst_1 = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc1)
|
| 18 |
+
%c-1_i64 = arith.constant -1 : i64 loc(#loc1)
|
| 19 |
+
%c2_i64 = arith.constant 2 : i64 loc(#loc1)
|
| 20 |
+
%cst_2 = arith.constant dense<true> : tensor<1024xi1> loc(#loc1)
|
| 21 |
+
%c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
|
| 22 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc81)
|
| 23 |
+
%xoffset_3 = arith.muli %xoffset, %c1024_i32 : i32 loc(#loc82)
|
| 24 |
+
%xindex = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc83)
|
| 25 |
+
%xindex_4 = tt.splat %xoffset_3 : i32 -> tensor<1024xi32> loc(#loc84)
|
| 26 |
+
%xindex_5 = arith.addi %xindex_4, %xindex : tensor<1024xi32> loc(#loc84)
|
| 27 |
+
%xmask = tt.splat %xnumel : i32 -> tensor<1024xi32> loc(#loc85)
|
| 28 |
+
%xmask_6 = arith.cmpi slt, %xindex_5, %xmask : tensor<1024xi32> loc(#loc85)
|
| 29 |
+
%x2 = arith.extsi %xindex_5 : tensor<1024xi32> to tensor<1024xi64> loc(#loc86)
|
| 30 |
+
%x2_7 = tt.splat %ks0 : i64 -> tensor<1024xi64> loc(#loc86)
|
| 31 |
+
%x2_8 = arith.divsi %x2, %x2_7 : tensor<1024xi64> loc(#loc86)
|
| 32 |
+
%x2_9 = tt.splat %ks1 : i64 -> tensor<1024xi64> loc(#loc87)
|
| 33 |
+
%x2_10 = arith.remsi %x2_8, %x2_9 : tensor<1024xi64> loc(#loc87)
|
| 34 |
+
%x0 = tt.splat %ks3 : i64 -> tensor<1024xi64> loc(#loc88)
|
| 35 |
+
%x0_11 = arith.remsi %x2, %x0 : tensor<1024xi64> loc(#loc88)
|
| 36 |
+
%x5 = arith.divsi %x2, %x0 : tensor<1024xi64> loc(#loc89)
|
| 37 |
+
%tmp0 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc90)
|
| 38 |
+
%tmp0_12 = tt.addptr %tmp0, %xindex_5 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc90)
|
| 39 |
+
%tmp0_13 = tt.load %tmp0_12, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc91)
|
| 40 |
+
%tmp0_14 = arith.extf %tmp0_13 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc92)
|
| 41 |
+
%tmp1 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1024x!tt.ptr<i64>> loc(#loc93)
|
| 42 |
+
%tmp1_15 = tt.addptr %tmp1, %x2_10 : tensor<1024x!tt.ptr<i64>>, tensor<1024xi64> loc(#loc93)
|
| 43 |
+
%tmp1_16 = tt.load %tmp1_15, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<i64>> loc(#loc94)
|
| 44 |
+
%tmp3 = tt.splat %ks2 : i64 -> tensor<1024xi64> loc(#loc95)
|
| 45 |
+
%tmp3_17 = arith.addi %tmp1_16, %tmp3 : tensor<1024xi64> loc(#loc95)
|
| 46 |
+
%tmp4 = arith.cmpi slt, %tmp1_16, %cst_0 : tensor<1024xi64> loc(#loc96)
|
| 47 |
+
%tmp5 = arith.select %tmp4, %tmp3_17, %tmp1_16 : tensor<1024xi1>, tensor<1024xi64> loc(#loc97)
|
| 48 |
+
%0 = arith.cmpi sge, %tmp5, %cst_0 : tensor<1024xi64> loc(#loc19)
|
| 49 |
+
%1 = arith.cmpi slt, %tmp5, %tmp3 : tensor<1024xi64> loc(#loc20)
|
| 50 |
+
%2 = arith.andi %0, %1 : tensor<1024xi1> loc(#loc21)
|
| 51 |
+
%3 = arith.xori %xmask_6, %cst_2 : tensor<1024xi1> loc(#loc22)
|
| 52 |
+
%4 = arith.ori %2, %3 : tensor<1024xi1> loc(#loc23)
|
| 53 |
+
tt.assert %4, "index out of bounds: 0 <= tmp5 < ks2" : tensor<1024xi1> loc(#loc24)
|
| 54 |
+
%tmp7 = arith.muli %x0, %tmp5 : tensor<1024xi64> loc(#loc98)
|
| 55 |
+
%tmp7_18 = arith.addi %x0_11, %tmp7 : tensor<1024xi64> loc(#loc99)
|
| 56 |
+
%tmp7_19 = tt.splat %in_ptr2 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc100)
|
| 57 |
+
%tmp7_20 = tt.addptr %tmp7_19, %tmp7_18 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc100)
|
| 58 |
+
%tmp7_21 = tt.load %tmp7_20, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc101)
|
| 59 |
+
%tmp7_22 = arith.extf %tmp7_21 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc102)
|
| 60 |
+
%tmp8 = arith.mulf %tmp0_14, %tmp7_22 : tensor<1024xf32> loc(#loc103)
|
| 61 |
+
%tmp12 = arith.divsi %ks3, %c2_i64 : i64 loc(#loc104)
|
| 62 |
+
%tmp12_23 = arith.subi %ks3, %tmp12 : i64 loc(#loc105)
|
| 63 |
+
%tmp13 = tt.splat %tmp12_23 : i64 -> tensor<1024xi64> loc(#loc106)
|
| 64 |
+
%tmp13_24 = arith.cmpi slt, %x0_11, %tmp13 : tensor<1024xi64> loc(#loc106)
|
| 65 |
+
%tmp14 = arith.muli %x0, %x5 : tensor<1024xi64> loc(#loc107)
|
| 66 |
+
%tmp14_25 = tt.splat %tmp12 : i64 -> tensor<1024xi64> loc(#loc108)
|
| 67 |
+
%tmp14_26 = arith.addi %tmp14, %tmp14_25 : tensor<1024xi64> loc(#loc108)
|
| 68 |
+
%tmp14_27 = arith.addi %tmp14_26, %x0_11 : tensor<1024xi64> loc(#loc109)
|
| 69 |
+
%tmp14_28 = tt.addptr %tmp0, %tmp14_27 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc110)
|
| 70 |
+
%tmp14_29 = arith.andi %tmp13_24, %xmask_6 : tensor<1024xi1> loc(#loc111)
|
| 71 |
+
%tmp14_30 = tt.load %tmp14_28, %tmp14_29, %cst evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc112)
|
| 72 |
+
%tmp14_31 = arith.extf %tmp14_30 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc113)
|
| 73 |
+
%tmp15 = arith.subf %cst_1, %tmp14_31 : tensor<1024xf32> loc(#loc114)
|
| 74 |
+
%tmp18 = arith.cmpi sge, %x0_11, %tmp13 : tensor<1024xi64> loc(#loc115)
|
| 75 |
+
%tmp21 = arith.muli %ks3, %c-1_i64 : i64 loc(#loc116)
|
| 76 |
+
%tmp21_32 = tt.splat %tmp21 : i64 -> tensor<1024xi64> loc(#loc117)
|
| 77 |
+
%tmp21_33 = arith.addi %x0_11, %tmp21_32 : tensor<1024xi64> loc(#loc117)
|
| 78 |
+
%tmp21_34 = arith.addi %tmp21_33, %tmp14_25 : tensor<1024xi64> loc(#loc118)
|
| 79 |
+
%tmp21_35 = arith.addi %tmp14, %tmp21_34 : tensor<1024xi64> loc(#loc119)
|
| 80 |
+
%tmp21_36 = tt.addptr %tmp0, %tmp21_35 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc120)
|
| 81 |
+
%tmp21_37 = arith.andi %tmp18, %xmask_6 : tensor<1024xi1> loc(#loc121)
|
| 82 |
+
%tmp21_38 = tt.load %tmp21_36, %tmp21_37, %cst evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc122)
|
| 83 |
+
%tmp21_39 = arith.extf %tmp21_38 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc123)
|
| 84 |
+
%tmp22 = arith.select %tmp13_24, %tmp15, %tmp21_39 : tensor<1024xi1>, tensor<1024xf32> loc(#loc135)
|
| 85 |
+
%tmp24 = tt.splat %ks4 : i64 -> tensor<1024xi64> loc(#loc126)
|
| 86 |
+
%tmp24_40 = arith.addi %tmp1_16, %tmp24 : tensor<1024xi64> loc(#loc126)
|
| 87 |
+
%tmp25 = arith.select %tmp4, %tmp24_40, %tmp1_16 : tensor<1024xi1>, tensor<1024xi64> loc(#loc127)
|
| 88 |
+
%5 = arith.cmpi sge, %tmp25, %cst_0 : tensor<1024xi64> loc(#loc55)
|
| 89 |
+
%6 = arith.cmpi slt, %tmp25, %tmp24 : tensor<1024xi64> loc(#loc56)
|
| 90 |
+
%7 = arith.andi %5, %6 : tensor<1024xi1> loc(#loc57)
|
| 91 |
+
%8 = arith.ori %7, %3 : tensor<1024xi1> loc(#loc58)
|
| 92 |
+
tt.assert %8, "index out of bounds: 0 <= tmp25 < ks4" : tensor<1024xi1> loc(#loc59)
|
| 93 |
+
%tmp27 = arith.muli %x0, %tmp25 : tensor<1024xi64> loc(#loc128)
|
| 94 |
+
%tmp27_41 = arith.addi %x0_11, %tmp27 : tensor<1024xi64> loc(#loc129)
|
| 95 |
+
%tmp27_42 = tt.splat %in_ptr3 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc130)
|
| 96 |
+
%tmp27_43 = tt.addptr %tmp27_42, %tmp27_41 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi64> loc(#loc130)
|
| 97 |
+
%tmp27_44 = tt.load %tmp27_43, %xmask_6 evictionPolicy = evict_last : tensor<1024x!tt.ptr<bf16>> loc(#loc131)
|
| 98 |
+
%tmp27_45 = arith.extf %tmp27_44 : tensor<1024xbf16> to tensor<1024xf32> loc(#loc132)
|
| 99 |
+
%tmp28 = arith.mulf %tmp22, %tmp27_45 : tensor<1024xf32> loc(#loc133)
|
| 100 |
+
%tmp29 = arith.addf %tmp8, %tmp28 : tensor<1024xf32> loc(#loc134)
|
| 101 |
+
%9 = tt.splat %out_ptr0 : !tt.ptr<bf16> -> tensor<1024x!tt.ptr<bf16>> loc(#loc67)
|
| 102 |
+
%10 = tt.addptr %9, %xindex_5 : tensor<1024x!tt.ptr<bf16>>, tensor<1024xi32> loc(#loc67)
|
| 103 |
+
%11 = arith.truncf %tmp29 : tensor<1024xf32> to tensor<1024xbf16> loc(#loc68)
|
| 104 |
+
tt.store %10, %11, %xmask_6 : tensor<1024x!tt.ptr<bf16>> loc(#loc68)
|
| 105 |
+
tt.return loc(#loc69)
|
| 106 |
+
} loc(#loc)
|
| 107 |
+
} loc(#loc)
|
| 108 |
+
#loc1 = loc(unknown)
|
| 109 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":19:28)
|
| 110 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":19:33)
|
| 111 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":20:36)
|
| 112 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":20:23)
|
| 113 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":21:21)
|
| 114 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":23:21)
|
| 115 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":23:28)
|
| 116 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":24:19)
|
| 117 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":25:19)
|
| 118 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:30)
|
| 119 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:35)
|
| 120 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":26:75)
|
| 121 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":27:30)
|
| 122 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":27:35)
|
| 123 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":29:18)
|
| 124 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":30:18)
|
| 125 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":31:32)
|
| 126 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:28)
|
| 127 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:44)
|
| 128 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:37)
|
| 129 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:54)
|
| 130 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:52)
|
| 131 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":32:62)
|
| 132 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:39)
|
| 133 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:35)
|
| 134 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:30)
|
| 135 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:46)
|
| 136 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":33:86)
|
| 137 |
+
#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":34:18)
|
| 138 |
+
#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:31)
|
| 139 |
+
#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":38:18)
|
| 140 |
+
#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":39:19)
|
| 141 |
+
#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:35)
|
| 142 |
+
#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:41)
|
| 143 |
+
#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:54)
|
| 144 |
+
#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:31)
|
| 145 |
+
#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:68)
|
| 146 |
+
#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:60)
|
| 147 |
+
#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":40:119)
|
| 148 |
+
#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":41:13)
|
| 149 |
+
#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":44:20)
|
| 150 |
+
#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:52)
|
| 151 |
+
#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:47)
|
| 152 |
+
#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:60)
|
| 153 |
+
#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:41)
|
| 154 |
+
#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:31)
|
| 155 |
+
#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:81)
|
| 156 |
+
#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:73)
|
| 157 |
+
#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":47:132)
|
| 158 |
+
#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":48:35)
|
| 159 |
+
#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":43:35)
|
| 160 |
+
#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":50:19)
|
| 161 |
+
#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":51:34)
|
| 162 |
+
#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:28)
|
| 163 |
+
#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:46)
|
| 164 |
+
#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:38)
|
| 165 |
+
#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:54)
|
| 166 |
+
#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":52:64)
|
| 167 |
+
#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:40)
|
| 168 |
+
#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:36)
|
| 169 |
+
#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:31)
|
| 170 |
+
#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:48)
|
| 171 |
+
#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":53:88)
|
| 172 |
+
#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":54:20)
|
| 173 |
+
#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":55:19)
|
| 174 |
+
#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:25)
|
| 175 |
+
#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:37)
|
| 176 |
+
#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/rv/crvxbztkr372fgdn7bgrud22s3wmd2isidwo4ek4hldn56tuv2dj.py":56:4)
|
| 177 |
+
#loc81 = loc("xoffset"(#loc2))
|
| 178 |
+
#loc82 = loc("xoffset"(#loc3))
|
| 179 |
+
#loc83 = loc("xindex"(#loc4))
|
| 180 |
+
#loc84 = loc("xindex"(#loc5))
|
| 181 |
+
#loc85 = loc("xmask"(#loc6))
|
| 182 |
+
#loc86 = loc("x2"(#loc7))
|
| 183 |
+
#loc87 = loc("x2"(#loc8))
|
| 184 |
+
#loc88 = loc("x0"(#loc9))
|
| 185 |
+
#loc89 = loc("x5"(#loc10))
|
| 186 |
+
#loc90 = loc("tmp0"(#loc11))
|
| 187 |
+
#loc91 = loc("tmp0"(#loc12))
|
| 188 |
+
#loc92 = loc("tmp0"(#loc13))
|
| 189 |
+
#loc93 = loc("tmp1"(#loc14))
|
| 190 |
+
#loc94 = loc("tmp1"(#loc15))
|
| 191 |
+
#loc95 = loc("tmp3"(#loc16))
|
| 192 |
+
#loc96 = loc("tmp4"(#loc17))
|
| 193 |
+
#loc97 = loc("tmp5"(#loc18))
|
| 194 |
+
#loc98 = loc("tmp7"(#loc25))
|
| 195 |
+
#loc99 = loc("tmp7"(#loc26))
|
| 196 |
+
#loc100 = loc("tmp7"(#loc27))
|
| 197 |
+
#loc101 = loc("tmp7"(#loc28))
|
| 198 |
+
#loc102 = loc("tmp7"(#loc29))
|
| 199 |
+
#loc103 = loc("tmp8"(#loc30))
|
| 200 |
+
#loc104 = loc("tmp12"(#loc31))
|
| 201 |
+
#loc105 = loc("tmp12"(#loc32))
|
| 202 |
+
#loc106 = loc("tmp13"(#loc33))
|
| 203 |
+
#loc107 = loc("tmp14"(#loc34))
|
| 204 |
+
#loc108 = loc("tmp14"(#loc35))
|
| 205 |
+
#loc109 = loc("tmp14"(#loc36))
|
| 206 |
+
#loc110 = loc("tmp14"(#loc37))
|
| 207 |
+
#loc111 = loc("tmp14"(#loc38))
|
| 208 |
+
#loc112 = loc("tmp14"(#loc39))
|
| 209 |
+
#loc113 = loc("tmp14"(#loc40))
|
| 210 |
+
#loc114 = loc("tmp15"(#loc41))
|
| 211 |
+
#loc115 = loc("tmp18"(#loc42))
|
| 212 |
+
#loc116 = loc("tmp21"(#loc43))
|
| 213 |
+
#loc117 = loc("tmp21"(#loc44))
|
| 214 |
+
#loc118 = loc("tmp21"(#loc45))
|
| 215 |
+
#loc119 = loc("tmp21"(#loc46))
|
| 216 |
+
#loc120 = loc("tmp21"(#loc47))
|
| 217 |
+
#loc121 = loc("tmp21"(#loc48))
|
| 218 |
+
#loc122 = loc("tmp21"(#loc49))
|
| 219 |
+
#loc123 = loc("tmp21"(#loc50))
|
| 220 |
+
#loc124 = loc("tmp22"(#loc51))
|
| 221 |
+
#loc125 = loc("tmp17"(#loc52))
|
| 222 |
+
#loc126 = loc("tmp24"(#loc53))
|
| 223 |
+
#loc127 = loc("tmp25"(#loc54))
|
| 224 |
+
#loc128 = loc("tmp27"(#loc60))
|
| 225 |
+
#loc129 = loc("tmp27"(#loc61))
|
| 226 |
+
#loc130 = loc("tmp27"(#loc62))
|
| 227 |
+
#loc131 = loc("tmp27"(#loc63))
|
| 228 |
+
#loc132 = loc("tmp27"(#loc64))
|
| 229 |
+
#loc133 = loc("tmp28"(#loc65))
|
| 230 |
+
#loc134 = loc("tmp29"(#loc66))
|
| 231 |
+
#loc135 = loc(fused[#loc124, #loc125])
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/__grp__triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.source", "triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttir", "triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttgir", "triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.llir", "triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ptx", "triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.cubin", "triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"hash": "e06ef592ad45788917e8ca2aaf880ce2d9dc1b17df1f92a71a87d94ada7fc9a8", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 8, "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": 16384, "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_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2"}
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.llir
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ptx
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.source
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttgir
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4BXPLEVNIV4ISF7IZIVK7CAM4LM5YGYX34PZFJY2Q7MUVWT7ZGUA/triton_per_fused__to_copy_arange_bitwise_and_eq_gt_index_put_lt_new_zeros_scalar_tensor_sort_sum_unsqueeze_view_where_2.ttir
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/__grp__triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/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/1/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/1/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/1/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/1/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/1/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/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.cubin
ADDED
|
Binary file (86.4 kB). View file
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 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/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.llir
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ptx
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.source
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UGGVGDRILN3TLIINH7RHS2I4LFNQRBLGHV63VTW6KCPABX4L4FA/triton_per_fused__to_copy_clone_slice_sort_sum_transpose_3.ttir
ADDED
|
@@ -0,0 +1,799 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":18:0)
|
| 2 |
+
#loc1 = loc(unknown)
|
| 3 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":41:67)
|
| 4 |
+
#loc23 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":662:12)
|
| 5 |
+
#loc28 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":634:73)
|
| 6 |
+
#loc32 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:51)
|
| 7 |
+
#loc37 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:53)
|
| 8 |
+
#loc46 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:50)
|
| 9 |
+
#loc51 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:51)
|
| 10 |
+
#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":45:26)
|
| 11 |
+
#loc80 = loc("in_ptr0"(#loc))
|
| 12 |
+
#loc81 = loc("out_ptr2"(#loc))
|
| 13 |
+
#loc82 = loc("out_ptr3"(#loc))
|
| 14 |
+
#loc83 = loc("xnumel"(#loc))
|
| 15 |
+
#loc84 = loc("r0_numel"(#loc))
|
| 16 |
+
#loc106 = loc(callsite(#loc23 at #loc2))
|
| 17 |
+
#loc113 = loc("ileft"(#loc32))
|
| 18 |
+
#loc117 = loc("iright"(#loc37))
|
| 19 |
+
#loc126 = loc("left_idx"(#loc46))
|
| 20 |
+
#loc131 = loc("right_idx"(#loc51))
|
| 21 |
+
#loc150 = loc("tmp11"(#loc70))
|
| 22 |
+
#loc157 = loc(callsite(#loc28 at #loc106))
|
| 23 |
+
#loc161 = loc(callsite(#loc1 at #loc150))
|
| 24 |
+
#loc165 = loc(callsite(#loc113 at #loc157))
|
| 25 |
+
#loc169 = loc(callsite(#loc117 at #loc157))
|
| 26 |
+
#loc177 = loc(callsite(#loc126 at #loc157))
|
| 27 |
+
#loc182 = loc(callsite(#loc131 at #loc157))
|
| 28 |
+
#loc202 = loc(callsite(#loc1 at #loc165))
|
| 29 |
+
#loc204 = loc(callsite(#loc1 at #loc169))
|
| 30 |
+
#loc207 = loc(callsite(#loc1 at #loc177))
|
| 31 |
+
#loc210 = loc(callsite(#loc1 at #loc182))
|
| 32 |
+
module {
|
| 33 |
+
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} {
|
| 34 |
+
%cst = arith.constant dense<1> : tensor<1x2x1xi32> loc(#loc85)
|
| 35 |
+
%cst_0 = arith.constant dense<0> : tensor<32x16xi32> loc(#loc1)
|
| 36 |
+
%tmp10 = arith.constant dense<0> : tensor<32x16xi64> loc(#loc86)
|
| 37 |
+
%tmp0 = arith.constant dense<272> : tensor<32x1xi32> loc(#loc87)
|
| 38 |
+
%tmp0_1 = arith.constant dense<17> : tensor<1x16xi32> loc(#loc88)
|
| 39 |
+
%cst_2 = arith.constant dense<16> : tensor<32x1xi32> loc(#loc1)
|
| 40 |
+
%xmask = arith.constant dense<128> : tensor<32x1xi32> loc(#loc89)
|
| 41 |
+
%c32_i32 = arith.constant 32 : i32 loc(#loc1)
|
| 42 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc90)
|
| 43 |
+
%xoffset_3 = arith.muli %xoffset, %c32_i32 : i32 loc(#loc91)
|
| 44 |
+
%xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc92)
|
| 45 |
+
%xindex_4 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc93)
|
| 46 |
+
%xindex_5 = tt.splat %xoffset_3 : i32 -> tensor<32x1xi32> loc(#loc94)
|
| 47 |
+
%xindex_6 = arith.addi %xindex_5, %xindex_4 : tensor<32x1xi32> loc(#loc94)
|
| 48 |
+
%xmask_7 = arith.cmpi slt, %xindex_6, %xmask : tensor<32x1xi32> loc(#loc89)
|
| 49 |
+
%r0_index = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc95)
|
| 50 |
+
%r0_index_8 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc96)
|
| 51 |
+
%x0 = arith.remsi %xindex_6, %cst_2 : tensor<32x1xi32> loc(#loc97)
|
| 52 |
+
%x1 = arith.divsi %xindex_6, %cst_2 : tensor<32x1xi32> loc(#loc98)
|
| 53 |
+
%tmp0_9 = arith.muli %r0_index_8, %tmp0_1 : tensor<1x16xi32> loc(#loc88)
|
| 54 |
+
%tmp0_10 = tt.broadcast %x0 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc99)
|
| 55 |
+
%tmp0_11 = tt.broadcast %tmp0_9 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc99)
|
| 56 |
+
%tmp0_12 = arith.addi %tmp0_10, %tmp0_11 : tensor<32x16xi32> loc(#loc99)
|
| 57 |
+
%tmp0_13 = arith.muli %x1, %tmp0 : tensor<32x1xi32> loc(#loc87)
|
| 58 |
+
%tmp0_14 = tt.broadcast %tmp0_13 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc100)
|
| 59 |
+
%tmp0_15 = arith.addi %tmp0_12, %tmp0_14 : tensor<32x16xi32> loc(#loc100)
|
| 60 |
+
%tmp0_16 = tt.splat %in_ptr0 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>> loc(#loc101)
|
| 61 |
+
%tmp0_17 = tt.addptr %tmp0_16, %tmp0_15 : tensor<32x16x!tt.ptr<i32>>, tensor<32x16xi32> loc(#loc101)
|
| 62 |
+
%tmp0_18 = tt.broadcast %xmask_7 : tensor<32x1xi1> -> tensor<32x16xi1> loc(#loc102)
|
| 63 |
+
%tmp0_19 = tt.load %tmp0_17, %tmp0_18, %cst_0 : tensor<32x16x!tt.ptr<i32>> loc(#loc102)
|
| 64 |
+
%tmp2 = arith.trunci %r0_index_8 : tensor<1x16xi32> to tensor<1x16xi16> loc(#loc103)
|
| 65 |
+
%tmp4 = tt.broadcast %tmp2 : tensor<1x16xi16> -> tensor<32x16xi16> loc(#loc104)
|
| 66 |
+
%flip = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32> loc(#loc153)
|
| 67 |
+
%flip_20 = tt.expand_dims %flip {axis = 0 : i32} : tensor<2xi32> -> tensor<1x2xi32> loc(#loc154)
|
| 68 |
+
%flip_21 = tt.expand_dims %flip_20 {axis = 2 : i32} : tensor<1x2xi32> -> tensor<1x2x1xi32> loc(#loc154)
|
| 69 |
+
%flip_22 = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<128x2x2xi32> loc(#loc155)
|
| 70 |
+
%flip_23 = tt.reshape %flip_22 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc156)
|
| 71 |
+
%y = tt.reshape %tmp0_19 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
|
| 72 |
+
%left_mask = arith.subi %cst, %flip_21 : tensor<1x2x1xi32> loc(#loc163)
|
| 73 |
+
%ileft = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<256x2x1xi32> loc(#loc164)
|
| 74 |
+
%ileft_24 = arith.muli %y, %ileft : tensor<256x2x1xi32> loc(#loc164)
|
| 75 |
+
%ileft_25 = "tt.reduce"(%ileft_24) <{axis = 1 : i32}> ({
|
| 76 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 77 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 78 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 79 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
|
| 80 |
+
%ileft_26 = tt.expand_dims %ileft_25 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
|
| 81 |
+
%ileft_27 = tt.broadcast %ileft_26 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
|
| 82 |
+
%iright = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<256x2x1xi32> loc(#loc168)
|
| 83 |
+
%iright_28 = arith.muli %y, %iright : tensor<256x2x1xi32> loc(#loc168)
|
| 84 |
+
%iright_29 = "tt.reduce"(%iright_28) <{axis = 1 : i32}> ({
|
| 85 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 86 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 87 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 88 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
|
| 89 |
+
%iright_30 = tt.expand_dims %iright_29 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
|
| 90 |
+
%iright_31 = tt.broadcast %iright_30 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
|
| 91 |
+
%ileft_32 = tt.reshape %ileft_27 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 92 |
+
%iright_33 = tt.reshape %iright_31 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 93 |
+
%y_idx = tt.reshape %tmp4 : tensor<32x16xi16> -> tensor<256x2x1xi16> loc(#loc174)
|
| 94 |
+
%left_idx = arith.trunci %left_mask : tensor<1x2x1xi32> to tensor<1x2x1xi16> loc(#loc175)
|
| 95 |
+
%left_idx_34 = tt.broadcast %left_idx : tensor<1x2x1xi16> -> tensor<256x2x1xi16> loc(#loc176)
|
| 96 |
+
%left_idx_35 = arith.muli %y_idx, %left_idx_34 : tensor<256x2x1xi16> loc(#loc176)
|
| 97 |
+
%input = arith.extsi %left_idx_35 : tensor<256x2x1xi16> to tensor<256x2x1xi32> loc(#loc205)
|
| 98 |
+
%left_idx_36 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
|
| 99 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 100 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 101 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 102 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
|
| 103 |
+
%left_idx_37 = tt.expand_dims %left_idx_36 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
|
| 104 |
+
%left_idx_38 = tt.broadcast %left_idx_37 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
|
| 105 |
+
%right_idx = arith.trunci %flip_21 : tensor<1x2x1xi32> to tensor<1x2x1xi16> loc(#loc180)
|
| 106 |
+
%right_idx_39 = tt.broadcast %right_idx : tensor<1x2x1xi16> -> tensor<256x2x1xi16> loc(#loc181)
|
| 107 |
+
%right_idx_40 = arith.muli %y_idx, %right_idx_39 : tensor<256x2x1xi16> loc(#loc181)
|
| 108 |
+
%input_41 = arith.extsi %right_idx_40 : tensor<256x2x1xi16> to tensor<256x2x1xi32> loc(#loc208)
|
| 109 |
+
%right_idx_42 = "tt.reduce"(%input_41) <{axis = 1 : i32}> ({
|
| 110 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 111 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 112 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 113 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
|
| 114 |
+
%right_idx_43 = tt.expand_dims %right_idx_42 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
|
| 115 |
+
%right_idx_44 = tt.broadcast %right_idx_43 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
|
| 116 |
+
%left_idx_45 = tt.reshape %left_idx_38 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 117 |
+
%right_idx_46 = tt.reshape %right_idx_44 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 118 |
+
%cond = arith.cmpi slt, %ileft_32, %iright_33 : tensor<32x16xi32> loc(#loc187)
|
| 119 |
+
%eq = arith.cmpi eq, %ileft_32, %iright_33 : tensor<32x16xi32> loc(#loc188)
|
| 120 |
+
%cond_47 = arith.cmpi sgt, %left_idx_45, %right_idx_46 : tensor<32x16xi32> loc(#loc189)
|
| 121 |
+
%cond_48 = arith.andi %eq, %cond_47 : tensor<32x16xi1> loc(#loc190)
|
| 122 |
+
%cond_49 = arith.ori %cond, %cond_48 : tensor<32x16xi1> loc(#loc191)
|
| 123 |
+
%cond_50 = arith.extui %cond_49 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
|
| 124 |
+
%cond_51 = arith.xori %cond_50, %flip_23 : tensor<32x16xi32> loc(#loc192)
|
| 125 |
+
%cond_52 = arith.cmpi ne, %cond_51, %cst_0 : tensor<32x16xi32> loc(#loc193)
|
| 126 |
+
%ret = arith.xori %ileft_32, %iright_33 : tensor<32x16xi32> loc(#loc194)
|
| 127 |
+
%ret_53 = arith.select %cond_52, %ret, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 128 |
+
%ret_54 = arith.xori %tmp0_19, %ret_53 : tensor<32x16xi32> loc(#loc196)
|
| 129 |
+
%new_idxs = arith.xori %left_idx_45, %right_idx_46 : tensor<32x16xi32> loc(#loc197)
|
| 130 |
+
%new_idxs_55 = arith.select %cond_52, %new_idxs, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 131 |
+
%new_idxs_56 = arith.extsi %tmp2 : tensor<1x16xi16> to tensor<1x16xi32> loc(#loc199)
|
| 132 |
+
%new_idxs_57 = tt.broadcast %new_idxs_56 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc199)
|
| 133 |
+
%new_idxs_58 = arith.xori %new_idxs_57, %new_idxs_55 : tensor<32x16xi32> loc(#loc199)
|
| 134 |
+
%flip_59 = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<64x2x4xi32> loc(#loc155)
|
| 135 |
+
%flip_60 = tt.reshape %flip_59 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc156)
|
| 136 |
+
%y_61 = tt.reshape %ret_54 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc162)
|
| 137 |
+
%ileft_62 = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<128x2x2xi32> loc(#loc164)
|
| 138 |
+
%ileft_63 = arith.muli %y_61, %ileft_62 : tensor<128x2x2xi32> loc(#loc164)
|
| 139 |
+
%ileft_64 = "tt.reduce"(%ileft_63) <{axis = 1 : i32}> ({
|
| 140 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 141 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 142 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 143 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc201)
|
| 144 |
+
%ileft_65 = tt.expand_dims %ileft_64 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc166)
|
| 145 |
+
%ileft_66 = tt.broadcast %ileft_65 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc167)
|
| 146 |
+
%iright_67 = arith.muli %y_61, %flip_22 : tensor<128x2x2xi32> loc(#loc168)
|
| 147 |
+
%iright_68 = "tt.reduce"(%iright_67) <{axis = 1 : i32}> ({
|
| 148 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 149 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 150 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 151 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc203)
|
| 152 |
+
%iright_69 = tt.expand_dims %iright_68 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc170)
|
| 153 |
+
%iright_70 = tt.broadcast %iright_69 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc171)
|
| 154 |
+
%ileft_71 = tt.reshape %ileft_66 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 155 |
+
%iright_72 = tt.reshape %iright_70 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 156 |
+
%y_idx_73 = tt.reshape %new_idxs_58 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc174)
|
| 157 |
+
%left_idx_74 = arith.muli %y_idx_73, %ileft_62 : tensor<128x2x2xi32> loc(#loc176)
|
| 158 |
+
%left_idx_75 = "tt.reduce"(%left_idx_74) <{axis = 1 : i32}> ({
|
| 159 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 160 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 161 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 162 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc206)
|
| 163 |
+
%left_idx_76 = tt.expand_dims %left_idx_75 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc178)
|
| 164 |
+
%left_idx_77 = tt.broadcast %left_idx_76 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc179)
|
| 165 |
+
%right_idx_78 = arith.muli %y_idx_73, %flip_22 : tensor<128x2x2xi32> loc(#loc181)
|
| 166 |
+
%right_idx_79 = "tt.reduce"(%right_idx_78) <{axis = 1 : i32}> ({
|
| 167 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 168 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 169 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 170 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc209)
|
| 171 |
+
%right_idx_80 = tt.expand_dims %right_idx_79 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc183)
|
| 172 |
+
%right_idx_81 = tt.broadcast %right_idx_80 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc184)
|
| 173 |
+
%left_idx_82 = tt.reshape %left_idx_77 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 174 |
+
%right_idx_83 = tt.reshape %right_idx_81 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 175 |
+
%cond_84 = arith.cmpi slt, %ileft_71, %iright_72 : tensor<32x16xi32> loc(#loc187)
|
| 176 |
+
%eq_85 = arith.cmpi eq, %ileft_71, %iright_72 : tensor<32x16xi32> loc(#loc188)
|
| 177 |
+
%cond_86 = arith.cmpi sgt, %left_idx_82, %right_idx_83 : tensor<32x16xi32> loc(#loc189)
|
| 178 |
+
%cond_87 = arith.andi %eq_85, %cond_86 : tensor<32x16xi1> loc(#loc190)
|
| 179 |
+
%cond_88 = arith.ori %cond_84, %cond_87 : tensor<32x16xi1> loc(#loc191)
|
| 180 |
+
%cond_89 = arith.extui %cond_88 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
|
| 181 |
+
%cond_90 = arith.xori %cond_89, %flip_60 : tensor<32x16xi32> loc(#loc192)
|
| 182 |
+
%cond_91 = arith.cmpi ne, %cond_90, %cst_0 : tensor<32x16xi32> loc(#loc193)
|
| 183 |
+
%ret_92 = arith.xori %ileft_71, %iright_72 : tensor<32x16xi32> loc(#loc194)
|
| 184 |
+
%ret_93 = arith.select %cond_91, %ret_92, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 185 |
+
%ret_94 = arith.xori %ret_54, %ret_93 : tensor<32x16xi32> loc(#loc196)
|
| 186 |
+
%new_idxs_95 = arith.xori %left_idx_82, %right_idx_83 : tensor<32x16xi32> loc(#loc197)
|
| 187 |
+
%new_idxs_96 = arith.select %cond_91, %new_idxs_95, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 188 |
+
%new_idxs_97 = arith.xori %new_idxs_58, %new_idxs_96 : tensor<32x16xi32> loc(#loc199)
|
| 189 |
+
%y_98 = tt.reshape %ret_94 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
|
| 190 |
+
%ileft_99 = arith.muli %y_98, %ileft : tensor<256x2x1xi32> loc(#loc164)
|
| 191 |
+
%ileft_100 = "tt.reduce"(%ileft_99) <{axis = 1 : i32}> ({
|
| 192 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 193 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 194 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 195 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
|
| 196 |
+
%ileft_101 = tt.expand_dims %ileft_100 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
|
| 197 |
+
%ileft_102 = tt.broadcast %ileft_101 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
|
| 198 |
+
%iright_103 = arith.muli %y_98, %iright : tensor<256x2x1xi32> loc(#loc168)
|
| 199 |
+
%iright_104 = "tt.reduce"(%iright_103) <{axis = 1 : i32}> ({
|
| 200 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 201 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 202 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 203 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
|
| 204 |
+
%iright_105 = tt.expand_dims %iright_104 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
|
| 205 |
+
%iright_106 = tt.broadcast %iright_105 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
|
| 206 |
+
%ileft_107 = tt.reshape %ileft_102 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 207 |
+
%iright_108 = tt.reshape %iright_106 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 208 |
+
%y_idx_109 = tt.reshape %new_idxs_97 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc174)
|
| 209 |
+
%left_idx_110 = arith.muli %y_idx_109, %ileft : tensor<256x2x1xi32> loc(#loc176)
|
| 210 |
+
%left_idx_111 = "tt.reduce"(%left_idx_110) <{axis = 1 : i32}> ({
|
| 211 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 212 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 213 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 214 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
|
| 215 |
+
%left_idx_112 = tt.expand_dims %left_idx_111 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
|
| 216 |
+
%left_idx_113 = tt.broadcast %left_idx_112 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
|
| 217 |
+
%right_idx_114 = arith.muli %y_idx_109, %iright : tensor<256x2x1xi32> loc(#loc181)
|
| 218 |
+
%right_idx_115 = "tt.reduce"(%right_idx_114) <{axis = 1 : i32}> ({
|
| 219 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 220 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 221 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 222 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
|
| 223 |
+
%right_idx_116 = tt.expand_dims %right_idx_115 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
|
| 224 |
+
%right_idx_117 = tt.broadcast %right_idx_116 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
|
| 225 |
+
%left_idx_118 = tt.reshape %left_idx_113 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 226 |
+
%right_idx_119 = tt.reshape %right_idx_117 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 227 |
+
%cond_120 = arith.cmpi slt, %ileft_107, %iright_108 : tensor<32x16xi32> loc(#loc187)
|
| 228 |
+
%eq_121 = arith.cmpi eq, %ileft_107, %iright_108 : tensor<32x16xi32> loc(#loc188)
|
| 229 |
+
%cond_122 = arith.cmpi sgt, %left_idx_118, %right_idx_119 : tensor<32x16xi32> loc(#loc189)
|
| 230 |
+
%cond_123 = arith.andi %eq_121, %cond_122 : tensor<32x16xi1> loc(#loc190)
|
| 231 |
+
%cond_124 = arith.ori %cond_120, %cond_123 : tensor<32x16xi1> loc(#loc191)
|
| 232 |
+
%cond_125 = arith.extui %cond_124 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
|
| 233 |
+
%cond_126 = arith.xori %cond_125, %flip_60 : tensor<32x16xi32> loc(#loc192)
|
| 234 |
+
%cond_127 = arith.cmpi ne, %cond_126, %cst_0 : tensor<32x16xi32> loc(#loc193)
|
| 235 |
+
%ret_128 = arith.xori %ileft_107, %iright_108 : tensor<32x16xi32> loc(#loc194)
|
| 236 |
+
%ret_129 = arith.select %cond_127, %ret_128, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 237 |
+
%ret_130 = arith.xori %ret_94, %ret_129 : tensor<32x16xi32> loc(#loc196)
|
| 238 |
+
%new_idxs_131 = arith.xori %left_idx_118, %right_idx_119 : tensor<32x16xi32> loc(#loc197)
|
| 239 |
+
%new_idxs_132 = arith.select %cond_127, %new_idxs_131, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 240 |
+
%new_idxs_133 = arith.xori %new_idxs_97, %new_idxs_132 : tensor<32x16xi32> loc(#loc199)
|
| 241 |
+
%flip_134 = tt.broadcast %flip_21 : tensor<1x2x1xi32> -> tensor<32x2x8xi32> loc(#loc155)
|
| 242 |
+
%flip_135 = tt.reshape %flip_134 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc156)
|
| 243 |
+
%y_136 = tt.reshape %ret_130 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc162)
|
| 244 |
+
%ileft_137 = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<64x2x4xi32> loc(#loc164)
|
| 245 |
+
%ileft_138 = arith.muli %y_136, %ileft_137 : tensor<64x2x4xi32> loc(#loc164)
|
| 246 |
+
%ileft_139 = "tt.reduce"(%ileft_138) <{axis = 1 : i32}> ({
|
| 247 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 248 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 249 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 250 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc201)
|
| 251 |
+
%ileft_140 = tt.expand_dims %ileft_139 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc166)
|
| 252 |
+
%ileft_141 = tt.broadcast %ileft_140 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc167)
|
| 253 |
+
%iright_142 = arith.muli %y_136, %flip_59 : tensor<64x2x4xi32> loc(#loc168)
|
| 254 |
+
%iright_143 = "tt.reduce"(%iright_142) <{axis = 1 : i32}> ({
|
| 255 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 256 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 257 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 258 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc203)
|
| 259 |
+
%iright_144 = tt.expand_dims %iright_143 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc170)
|
| 260 |
+
%iright_145 = tt.broadcast %iright_144 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc171)
|
| 261 |
+
%ileft_146 = tt.reshape %ileft_141 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 262 |
+
%iright_147 = tt.reshape %iright_145 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 263 |
+
%y_idx_148 = tt.reshape %new_idxs_133 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc174)
|
| 264 |
+
%left_idx_149 = arith.muli %y_idx_148, %ileft_137 : tensor<64x2x4xi32> loc(#loc176)
|
| 265 |
+
%left_idx_150 = "tt.reduce"(%left_idx_149) <{axis = 1 : i32}> ({
|
| 266 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 267 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 268 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 269 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc206)
|
| 270 |
+
%left_idx_151 = tt.expand_dims %left_idx_150 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc178)
|
| 271 |
+
%left_idx_152 = tt.broadcast %left_idx_151 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc179)
|
| 272 |
+
%right_idx_153 = arith.muli %y_idx_148, %flip_59 : tensor<64x2x4xi32> loc(#loc181)
|
| 273 |
+
%right_idx_154 = "tt.reduce"(%right_idx_153) <{axis = 1 : i32}> ({
|
| 274 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 275 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 276 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 277 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc209)
|
| 278 |
+
%right_idx_155 = tt.expand_dims %right_idx_154 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc183)
|
| 279 |
+
%right_idx_156 = tt.broadcast %right_idx_155 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc184)
|
| 280 |
+
%left_idx_157 = tt.reshape %left_idx_152 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 281 |
+
%right_idx_158 = tt.reshape %right_idx_156 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 282 |
+
%cond_159 = arith.cmpi slt, %ileft_146, %iright_147 : tensor<32x16xi32> loc(#loc187)
|
| 283 |
+
%eq_160 = arith.cmpi eq, %ileft_146, %iright_147 : tensor<32x16xi32> loc(#loc188)
|
| 284 |
+
%cond_161 = arith.cmpi sgt, %left_idx_157, %right_idx_158 : tensor<32x16xi32> loc(#loc189)
|
| 285 |
+
%cond_162 = arith.andi %eq_160, %cond_161 : tensor<32x16xi1> loc(#loc190)
|
| 286 |
+
%cond_163 = arith.ori %cond_159, %cond_162 : tensor<32x16xi1> loc(#loc191)
|
| 287 |
+
%cond_164 = arith.extui %cond_163 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
|
| 288 |
+
%cond_165 = arith.xori %cond_164, %flip_135 : tensor<32x16xi32> loc(#loc192)
|
| 289 |
+
%cond_166 = arith.cmpi ne, %cond_165, %cst_0 : tensor<32x16xi32> loc(#loc193)
|
| 290 |
+
%ret_167 = arith.xori %ileft_146, %iright_147 : tensor<32x16xi32> loc(#loc194)
|
| 291 |
+
%ret_168 = arith.select %cond_166, %ret_167, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 292 |
+
%ret_169 = arith.xori %ret_130, %ret_168 : tensor<32x16xi32> loc(#loc196)
|
| 293 |
+
%new_idxs_170 = arith.xori %left_idx_157, %right_idx_158 : tensor<32x16xi32> loc(#loc197)
|
| 294 |
+
%new_idxs_171 = arith.select %cond_166, %new_idxs_170, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 295 |
+
%new_idxs_172 = arith.xori %new_idxs_133, %new_idxs_171 : tensor<32x16xi32> loc(#loc199)
|
| 296 |
+
%y_173 = tt.reshape %ret_169 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc162)
|
| 297 |
+
%ileft_174 = arith.muli %y_173, %ileft_62 : tensor<128x2x2xi32> loc(#loc164)
|
| 298 |
+
%ileft_175 = "tt.reduce"(%ileft_174) <{axis = 1 : i32}> ({
|
| 299 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 300 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 301 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 302 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc201)
|
| 303 |
+
%ileft_176 = tt.expand_dims %ileft_175 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc166)
|
| 304 |
+
%ileft_177 = tt.broadcast %ileft_176 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc167)
|
| 305 |
+
%iright_178 = arith.muli %y_173, %flip_22 : tensor<128x2x2xi32> loc(#loc168)
|
| 306 |
+
%iright_179 = "tt.reduce"(%iright_178) <{axis = 1 : i32}> ({
|
| 307 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 308 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 309 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 310 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc203)
|
| 311 |
+
%iright_180 = tt.expand_dims %iright_179 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc170)
|
| 312 |
+
%iright_181 = tt.broadcast %iright_180 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc171)
|
| 313 |
+
%ileft_182 = tt.reshape %ileft_177 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 314 |
+
%iright_183 = tt.reshape %iright_181 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 315 |
+
%y_idx_184 = tt.reshape %new_idxs_172 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc174)
|
| 316 |
+
%left_idx_185 = arith.muli %y_idx_184, %ileft_62 : tensor<128x2x2xi32> loc(#loc176)
|
| 317 |
+
%left_idx_186 = "tt.reduce"(%left_idx_185) <{axis = 1 : i32}> ({
|
| 318 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 319 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 320 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 321 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc206)
|
| 322 |
+
%left_idx_187 = tt.expand_dims %left_idx_186 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc178)
|
| 323 |
+
%left_idx_188 = tt.broadcast %left_idx_187 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc179)
|
| 324 |
+
%right_idx_189 = arith.muli %y_idx_184, %flip_22 : tensor<128x2x2xi32> loc(#loc181)
|
| 325 |
+
%right_idx_190 = "tt.reduce"(%right_idx_189) <{axis = 1 : i32}> ({
|
| 326 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 327 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 328 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 329 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc209)
|
| 330 |
+
%right_idx_191 = tt.expand_dims %right_idx_190 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc183)
|
| 331 |
+
%right_idx_192 = tt.broadcast %right_idx_191 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc184)
|
| 332 |
+
%left_idx_193 = tt.reshape %left_idx_188 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 333 |
+
%right_idx_194 = tt.reshape %right_idx_192 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 334 |
+
%cond_195 = arith.cmpi slt, %ileft_182, %iright_183 : tensor<32x16xi32> loc(#loc187)
|
| 335 |
+
%eq_196 = arith.cmpi eq, %ileft_182, %iright_183 : tensor<32x16xi32> loc(#loc188)
|
| 336 |
+
%cond_197 = arith.cmpi sgt, %left_idx_193, %right_idx_194 : tensor<32x16xi32> loc(#loc189)
|
| 337 |
+
%cond_198 = arith.andi %eq_196, %cond_197 : tensor<32x16xi1> loc(#loc190)
|
| 338 |
+
%cond_199 = arith.ori %cond_195, %cond_198 : tensor<32x16xi1> loc(#loc191)
|
| 339 |
+
%cond_200 = arith.extui %cond_199 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
|
| 340 |
+
%cond_201 = arith.xori %cond_200, %flip_135 : tensor<32x16xi32> loc(#loc192)
|
| 341 |
+
%cond_202 = arith.cmpi ne, %cond_201, %cst_0 : tensor<32x16xi32> loc(#loc193)
|
| 342 |
+
%ret_203 = arith.xori %ileft_182, %iright_183 : tensor<32x16xi32> loc(#loc194)
|
| 343 |
+
%ret_204 = arith.select %cond_202, %ret_203, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 344 |
+
%ret_205 = arith.xori %ret_169, %ret_204 : tensor<32x16xi32> loc(#loc196)
|
| 345 |
+
%new_idxs_206 = arith.xori %left_idx_193, %right_idx_194 : tensor<32x16xi32> loc(#loc197)
|
| 346 |
+
%new_idxs_207 = arith.select %cond_202, %new_idxs_206, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 347 |
+
%new_idxs_208 = arith.xori %new_idxs_172, %new_idxs_207 : tensor<32x16xi32> loc(#loc199)
|
| 348 |
+
%y_209 = tt.reshape %ret_205 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
|
| 349 |
+
%ileft_210 = arith.muli %y_209, %ileft : tensor<256x2x1xi32> loc(#loc164)
|
| 350 |
+
%ileft_211 = "tt.reduce"(%ileft_210) <{axis = 1 : i32}> ({
|
| 351 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 352 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 353 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 354 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
|
| 355 |
+
%ileft_212 = tt.expand_dims %ileft_211 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
|
| 356 |
+
%ileft_213 = tt.broadcast %ileft_212 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
|
| 357 |
+
%iright_214 = arith.muli %y_209, %iright : tensor<256x2x1xi32> loc(#loc168)
|
| 358 |
+
%iright_215 = "tt.reduce"(%iright_214) <{axis = 1 : i32}> ({
|
| 359 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 360 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 361 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 362 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
|
| 363 |
+
%iright_216 = tt.expand_dims %iright_215 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
|
| 364 |
+
%iright_217 = tt.broadcast %iright_216 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
|
| 365 |
+
%ileft_218 = tt.reshape %ileft_213 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 366 |
+
%iright_219 = tt.reshape %iright_217 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 367 |
+
%y_idx_220 = tt.reshape %new_idxs_208 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc174)
|
| 368 |
+
%left_idx_221 = arith.muli %y_idx_220, %ileft : tensor<256x2x1xi32> loc(#loc176)
|
| 369 |
+
%left_idx_222 = "tt.reduce"(%left_idx_221) <{axis = 1 : i32}> ({
|
| 370 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 371 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 372 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 373 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
|
| 374 |
+
%left_idx_223 = tt.expand_dims %left_idx_222 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
|
| 375 |
+
%left_idx_224 = tt.broadcast %left_idx_223 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
|
| 376 |
+
%right_idx_225 = arith.muli %y_idx_220, %iright : tensor<256x2x1xi32> loc(#loc181)
|
| 377 |
+
%right_idx_226 = "tt.reduce"(%right_idx_225) <{axis = 1 : i32}> ({
|
| 378 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 379 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 380 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 381 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
|
| 382 |
+
%right_idx_227 = tt.expand_dims %right_idx_226 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
|
| 383 |
+
%right_idx_228 = tt.broadcast %right_idx_227 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
|
| 384 |
+
%left_idx_229 = tt.reshape %left_idx_224 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 385 |
+
%right_idx_230 = tt.reshape %right_idx_228 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 386 |
+
%cond_231 = arith.cmpi slt, %ileft_218, %iright_219 : tensor<32x16xi32> loc(#loc187)
|
| 387 |
+
%eq_232 = arith.cmpi eq, %ileft_218, %iright_219 : tensor<32x16xi32> loc(#loc188)
|
| 388 |
+
%cond_233 = arith.cmpi sgt, %left_idx_229, %right_idx_230 : tensor<32x16xi32> loc(#loc189)
|
| 389 |
+
%cond_234 = arith.andi %eq_232, %cond_233 : tensor<32x16xi1> loc(#loc190)
|
| 390 |
+
%cond_235 = arith.ori %cond_231, %cond_234 : tensor<32x16xi1> loc(#loc191)
|
| 391 |
+
%cond_236 = arith.extui %cond_235 : tensor<32x16xi1> to tensor<32x16xi32> loc(#loc192)
|
| 392 |
+
%cond_237 = arith.xori %cond_236, %flip_135 : tensor<32x16xi32> loc(#loc192)
|
| 393 |
+
%cond_238 = arith.cmpi ne, %cond_237, %cst_0 : tensor<32x16xi32> loc(#loc193)
|
| 394 |
+
%ret_239 = arith.xori %ileft_218, %iright_219 : tensor<32x16xi32> loc(#loc194)
|
| 395 |
+
%ret_240 = arith.select %cond_238, %ret_239, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 396 |
+
%ret_241 = arith.xori %ret_205, %ret_240 : tensor<32x16xi32> loc(#loc196)
|
| 397 |
+
%new_idxs_242 = arith.xori %left_idx_229, %right_idx_230 : tensor<32x16xi32> loc(#loc197)
|
| 398 |
+
%new_idxs_243 = arith.select %cond_238, %new_idxs_242, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 399 |
+
%new_idxs_244 = arith.xori %new_idxs_208, %new_idxs_243 : tensor<32x16xi32> loc(#loc199)
|
| 400 |
+
%y_245 = tt.reshape %ret_241 : tensor<32x16xi32> -> tensor<32x2x8xi32> loc(#loc162)
|
| 401 |
+
%ileft_246 = tt.broadcast %left_mask : tensor<1x2x1xi32> -> tensor<32x2x8xi32> loc(#loc164)
|
| 402 |
+
%ileft_247 = arith.muli %y_245, %ileft_246 : tensor<32x2x8xi32> loc(#loc164)
|
| 403 |
+
%ileft_248 = "tt.reduce"(%ileft_247) <{axis = 1 : i32}> ({
|
| 404 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 405 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 406 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 407 |
+
}) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc201)
|
| 408 |
+
%ileft_249 = tt.expand_dims %ileft_248 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc166)
|
| 409 |
+
%ileft_250 = tt.broadcast %ileft_249 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc167)
|
| 410 |
+
%iright_251 = arith.muli %y_245, %flip_134 : tensor<32x2x8xi32> loc(#loc168)
|
| 411 |
+
%iright_252 = "tt.reduce"(%iright_251) <{axis = 1 : i32}> ({
|
| 412 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 413 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 414 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 415 |
+
}) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc203)
|
| 416 |
+
%iright_253 = tt.expand_dims %iright_252 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc170)
|
| 417 |
+
%iright_254 = tt.broadcast %iright_253 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc171)
|
| 418 |
+
%ileft_255 = tt.reshape %ileft_250 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 419 |
+
%iright_256 = tt.reshape %iright_254 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 420 |
+
%y_idx_257 = tt.reshape %new_idxs_244 : tensor<32x16xi32> -> tensor<32x2x8xi32> loc(#loc174)
|
| 421 |
+
%left_idx_258 = arith.muli %y_idx_257, %ileft_246 : tensor<32x2x8xi32> loc(#loc176)
|
| 422 |
+
%left_idx_259 = "tt.reduce"(%left_idx_258) <{axis = 1 : i32}> ({
|
| 423 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 424 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 425 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 426 |
+
}) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc206)
|
| 427 |
+
%left_idx_260 = tt.expand_dims %left_idx_259 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc178)
|
| 428 |
+
%left_idx_261 = tt.broadcast %left_idx_260 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc179)
|
| 429 |
+
%right_idx_262 = arith.muli %y_idx_257, %flip_134 : tensor<32x2x8xi32> loc(#loc181)
|
| 430 |
+
%right_idx_263 = "tt.reduce"(%right_idx_262) <{axis = 1 : i32}> ({
|
| 431 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 432 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 433 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 434 |
+
}) : (tensor<32x2x8xi32>) -> tensor<32x8xi32> loc(#loc209)
|
| 435 |
+
%right_idx_264 = tt.expand_dims %right_idx_263 {axis = 1 : i32} : tensor<32x8xi32> -> tensor<32x1x8xi32> loc(#loc183)
|
| 436 |
+
%right_idx_265 = tt.broadcast %right_idx_264 : tensor<32x1x8xi32> -> tensor<32x2x8xi32> loc(#loc184)
|
| 437 |
+
%left_idx_266 = tt.reshape %left_idx_261 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 438 |
+
%right_idx_267 = tt.reshape %right_idx_265 : tensor<32x2x8xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 439 |
+
%cond_268 = arith.cmpi slt, %ileft_255, %iright_256 : tensor<32x16xi32> loc(#loc187)
|
| 440 |
+
%eq_269 = arith.cmpi eq, %ileft_255, %iright_256 : tensor<32x16xi32> loc(#loc188)
|
| 441 |
+
%cond_270 = arith.cmpi sgt, %left_idx_266, %right_idx_267 : tensor<32x16xi32> loc(#loc189)
|
| 442 |
+
%cond_271 = arith.andi %eq_269, %cond_270 : tensor<32x16xi1> loc(#loc190)
|
| 443 |
+
%cond_272 = arith.ori %cond_268, %cond_271 : tensor<32x16xi1> loc(#loc191)
|
| 444 |
+
%ret_273 = arith.xori %ileft_255, %iright_256 : tensor<32x16xi32> loc(#loc194)
|
| 445 |
+
%ret_274 = arith.select %cond_272, %ret_273, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 446 |
+
%ret_275 = arith.xori %ret_241, %ret_274 : tensor<32x16xi32> loc(#loc196)
|
| 447 |
+
%new_idxs_276 = arith.xori %left_idx_266, %right_idx_267 : tensor<32x16xi32> loc(#loc197)
|
| 448 |
+
%new_idxs_277 = arith.select %cond_272, %new_idxs_276, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 449 |
+
%new_idxs_278 = arith.xori %new_idxs_244, %new_idxs_277 : tensor<32x16xi32> loc(#loc199)
|
| 450 |
+
%y_279 = tt.reshape %ret_275 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc162)
|
| 451 |
+
%ileft_280 = arith.muli %y_279, %ileft_137 : tensor<64x2x4xi32> loc(#loc164)
|
| 452 |
+
%ileft_281 = "tt.reduce"(%ileft_280) <{axis = 1 : i32}> ({
|
| 453 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 454 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 455 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 456 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc201)
|
| 457 |
+
%ileft_282 = tt.expand_dims %ileft_281 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc166)
|
| 458 |
+
%ileft_283 = tt.broadcast %ileft_282 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc167)
|
| 459 |
+
%iright_284 = arith.muli %y_279, %flip_59 : tensor<64x2x4xi32> loc(#loc168)
|
| 460 |
+
%iright_285 = "tt.reduce"(%iright_284) <{axis = 1 : i32}> ({
|
| 461 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 462 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 463 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 464 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc203)
|
| 465 |
+
%iright_286 = tt.expand_dims %iright_285 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc170)
|
| 466 |
+
%iright_287 = tt.broadcast %iright_286 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc171)
|
| 467 |
+
%ileft_288 = tt.reshape %ileft_283 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 468 |
+
%iright_289 = tt.reshape %iright_287 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 469 |
+
%y_idx_290 = tt.reshape %new_idxs_278 : tensor<32x16xi32> -> tensor<64x2x4xi32> loc(#loc174)
|
| 470 |
+
%left_idx_291 = arith.muli %y_idx_290, %ileft_137 : tensor<64x2x4xi32> loc(#loc176)
|
| 471 |
+
%left_idx_292 = "tt.reduce"(%left_idx_291) <{axis = 1 : i32}> ({
|
| 472 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 473 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 474 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 475 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc206)
|
| 476 |
+
%left_idx_293 = tt.expand_dims %left_idx_292 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc178)
|
| 477 |
+
%left_idx_294 = tt.broadcast %left_idx_293 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc179)
|
| 478 |
+
%right_idx_295 = arith.muli %y_idx_290, %flip_59 : tensor<64x2x4xi32> loc(#loc181)
|
| 479 |
+
%right_idx_296 = "tt.reduce"(%right_idx_295) <{axis = 1 : i32}> ({
|
| 480 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 481 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 482 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 483 |
+
}) : (tensor<64x2x4xi32>) -> tensor<64x4xi32> loc(#loc209)
|
| 484 |
+
%right_idx_297 = tt.expand_dims %right_idx_296 {axis = 1 : i32} : tensor<64x4xi32> -> tensor<64x1x4xi32> loc(#loc183)
|
| 485 |
+
%right_idx_298 = tt.broadcast %right_idx_297 : tensor<64x1x4xi32> -> tensor<64x2x4xi32> loc(#loc184)
|
| 486 |
+
%left_idx_299 = tt.reshape %left_idx_294 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 487 |
+
%right_idx_300 = tt.reshape %right_idx_298 : tensor<64x2x4xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 488 |
+
%cond_301 = arith.cmpi slt, %ileft_288, %iright_289 : tensor<32x16xi32> loc(#loc187)
|
| 489 |
+
%eq_302 = arith.cmpi eq, %ileft_288, %iright_289 : tensor<32x16xi32> loc(#loc188)
|
| 490 |
+
%cond_303 = arith.cmpi sgt, %left_idx_299, %right_idx_300 : tensor<32x16xi32> loc(#loc189)
|
| 491 |
+
%cond_304 = arith.andi %eq_302, %cond_303 : tensor<32x16xi1> loc(#loc190)
|
| 492 |
+
%cond_305 = arith.ori %cond_301, %cond_304 : tensor<32x16xi1> loc(#loc191)
|
| 493 |
+
%ret_306 = arith.xori %ileft_288, %iright_289 : tensor<32x16xi32> loc(#loc194)
|
| 494 |
+
%ret_307 = arith.select %cond_305, %ret_306, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 495 |
+
%ret_308 = arith.xori %ret_275, %ret_307 : tensor<32x16xi32> loc(#loc196)
|
| 496 |
+
%new_idxs_309 = arith.xori %left_idx_299, %right_idx_300 : tensor<32x16xi32> loc(#loc197)
|
| 497 |
+
%new_idxs_310 = arith.select %cond_305, %new_idxs_309, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 498 |
+
%new_idxs_311 = arith.xori %new_idxs_278, %new_idxs_310 : tensor<32x16xi32> loc(#loc199)
|
| 499 |
+
%y_312 = tt.reshape %ret_308 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc162)
|
| 500 |
+
%ileft_313 = arith.muli %y_312, %ileft_62 : tensor<128x2x2xi32> loc(#loc164)
|
| 501 |
+
%ileft_314 = "tt.reduce"(%ileft_313) <{axis = 1 : i32}> ({
|
| 502 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 503 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 504 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 505 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc201)
|
| 506 |
+
%ileft_315 = tt.expand_dims %ileft_314 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc166)
|
| 507 |
+
%ileft_316 = tt.broadcast %ileft_315 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc167)
|
| 508 |
+
%iright_317 = arith.muli %y_312, %flip_22 : tensor<128x2x2xi32> loc(#loc168)
|
| 509 |
+
%iright_318 = "tt.reduce"(%iright_317) <{axis = 1 : i32}> ({
|
| 510 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 511 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 512 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 513 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc203)
|
| 514 |
+
%iright_319 = tt.expand_dims %iright_318 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc170)
|
| 515 |
+
%iright_320 = tt.broadcast %iright_319 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc171)
|
| 516 |
+
%ileft_321 = tt.reshape %ileft_316 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 517 |
+
%iright_322 = tt.reshape %iright_320 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 518 |
+
%y_idx_323 = tt.reshape %new_idxs_311 : tensor<32x16xi32> -> tensor<128x2x2xi32> loc(#loc174)
|
| 519 |
+
%left_idx_324 = arith.muli %y_idx_323, %ileft_62 : tensor<128x2x2xi32> loc(#loc176)
|
| 520 |
+
%left_idx_325 = "tt.reduce"(%left_idx_324) <{axis = 1 : i32}> ({
|
| 521 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 522 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 523 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 524 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc206)
|
| 525 |
+
%left_idx_326 = tt.expand_dims %left_idx_325 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc178)
|
| 526 |
+
%left_idx_327 = tt.broadcast %left_idx_326 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc179)
|
| 527 |
+
%right_idx_328 = arith.muli %y_idx_323, %flip_22 : tensor<128x2x2xi32> loc(#loc181)
|
| 528 |
+
%right_idx_329 = "tt.reduce"(%right_idx_328) <{axis = 1 : i32}> ({
|
| 529 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 530 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 531 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 532 |
+
}) : (tensor<128x2x2xi32>) -> tensor<128x2xi32> loc(#loc209)
|
| 533 |
+
%right_idx_330 = tt.expand_dims %right_idx_329 {axis = 1 : i32} : tensor<128x2xi32> -> tensor<128x1x2xi32> loc(#loc183)
|
| 534 |
+
%right_idx_331 = tt.broadcast %right_idx_330 : tensor<128x1x2xi32> -> tensor<128x2x2xi32> loc(#loc184)
|
| 535 |
+
%left_idx_332 = tt.reshape %left_idx_327 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 536 |
+
%right_idx_333 = tt.reshape %right_idx_331 : tensor<128x2x2xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 537 |
+
%cond_334 = arith.cmpi slt, %ileft_321, %iright_322 : tensor<32x16xi32> loc(#loc187)
|
| 538 |
+
%eq_335 = arith.cmpi eq, %ileft_321, %iright_322 : tensor<32x16xi32> loc(#loc188)
|
| 539 |
+
%cond_336 = arith.cmpi sgt, %left_idx_332, %right_idx_333 : tensor<32x16xi32> loc(#loc189)
|
| 540 |
+
%cond_337 = arith.andi %eq_335, %cond_336 : tensor<32x16xi1> loc(#loc190)
|
| 541 |
+
%cond_338 = arith.ori %cond_334, %cond_337 : tensor<32x16xi1> loc(#loc191)
|
| 542 |
+
%ret_339 = arith.xori %ileft_321, %iright_322 : tensor<32x16xi32> loc(#loc194)
|
| 543 |
+
%ret_340 = arith.select %cond_338, %ret_339, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc195)
|
| 544 |
+
%ret_341 = arith.xori %ret_308, %ret_340 : tensor<32x16xi32> loc(#loc196)
|
| 545 |
+
%new_idxs_342 = arith.xori %left_idx_332, %right_idx_333 : tensor<32x16xi32> loc(#loc197)
|
| 546 |
+
%new_idxs_343 = arith.select %cond_338, %new_idxs_342, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 547 |
+
%new_idxs_344 = arith.xori %new_idxs_311, %new_idxs_343 : tensor<32x16xi32> loc(#loc199)
|
| 548 |
+
%y_345 = tt.reshape %ret_341 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc162)
|
| 549 |
+
%ileft_346 = arith.muli %y_345, %ileft : tensor<256x2x1xi32> loc(#loc164)
|
| 550 |
+
%ileft_347 = "tt.reduce"(%ileft_346) <{axis = 1 : i32}> ({
|
| 551 |
+
^bb0(%ileft_377: i32 loc(callsite(#loc1 at #loc165)), %ileft_378: i32 loc(callsite(#loc1 at #loc165))):
|
| 552 |
+
%ileft_379 = arith.addi %ileft_377, %ileft_378 : i32 loc(#loc211)
|
| 553 |
+
tt.reduce.return %ileft_379 : i32 loc(#loc201)
|
| 554 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc201)
|
| 555 |
+
%ileft_348 = tt.expand_dims %ileft_347 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc166)
|
| 556 |
+
%ileft_349 = tt.broadcast %ileft_348 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc167)
|
| 557 |
+
%iright_350 = arith.muli %y_345, %iright : tensor<256x2x1xi32> loc(#loc168)
|
| 558 |
+
%iright_351 = "tt.reduce"(%iright_350) <{axis = 1 : i32}> ({
|
| 559 |
+
^bb0(%iright_377: i32 loc(callsite(#loc1 at #loc169)), %iright_378: i32 loc(callsite(#loc1 at #loc169))):
|
| 560 |
+
%iright_379 = arith.addi %iright_377, %iright_378 : i32 loc(#loc212)
|
| 561 |
+
tt.reduce.return %iright_379 : i32 loc(#loc203)
|
| 562 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc203)
|
| 563 |
+
%iright_352 = tt.expand_dims %iright_351 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc170)
|
| 564 |
+
%iright_353 = tt.broadcast %iright_352 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc171)
|
| 565 |
+
%ileft_354 = tt.reshape %ileft_349 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc172)
|
| 566 |
+
%iright_355 = tt.reshape %iright_353 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc173)
|
| 567 |
+
%y_idx_356 = tt.reshape %new_idxs_344 : tensor<32x16xi32> -> tensor<256x2x1xi32> loc(#loc174)
|
| 568 |
+
%left_idx_357 = arith.muli %y_idx_356, %ileft : tensor<256x2x1xi32> loc(#loc176)
|
| 569 |
+
%left_idx_358 = "tt.reduce"(%left_idx_357) <{axis = 1 : i32}> ({
|
| 570 |
+
^bb0(%left_idx_377: i32 loc(callsite(#loc1 at #loc177)), %left_idx_378: i32 loc(callsite(#loc1 at #loc177))):
|
| 571 |
+
%left_idx_379 = arith.addi %left_idx_377, %left_idx_378 : i32 loc(#loc213)
|
| 572 |
+
tt.reduce.return %left_idx_379 : i32 loc(#loc206)
|
| 573 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc206)
|
| 574 |
+
%left_idx_359 = tt.expand_dims %left_idx_358 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc178)
|
| 575 |
+
%left_idx_360 = tt.broadcast %left_idx_359 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc179)
|
| 576 |
+
%right_idx_361 = arith.muli %y_idx_356, %iright : tensor<256x2x1xi32> loc(#loc181)
|
| 577 |
+
%right_idx_362 = "tt.reduce"(%right_idx_361) <{axis = 1 : i32}> ({
|
| 578 |
+
^bb0(%right_idx_377: i32 loc(callsite(#loc1 at #loc182)), %right_idx_378: i32 loc(callsite(#loc1 at #loc182))):
|
| 579 |
+
%right_idx_379 = arith.addi %right_idx_377, %right_idx_378 : i32 loc(#loc214)
|
| 580 |
+
tt.reduce.return %right_idx_379 : i32 loc(#loc209)
|
| 581 |
+
}) : (tensor<256x2x1xi32>) -> tensor<256x1xi32> loc(#loc209)
|
| 582 |
+
%right_idx_363 = tt.expand_dims %right_idx_362 {axis = 1 : i32} : tensor<256x1xi32> -> tensor<256x1x1xi32> loc(#loc183)
|
| 583 |
+
%right_idx_364 = tt.broadcast %right_idx_363 : tensor<256x1x1xi32> -> tensor<256x2x1xi32> loc(#loc184)
|
| 584 |
+
%left_idx_365 = tt.reshape %left_idx_360 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc185)
|
| 585 |
+
%right_idx_366 = tt.reshape %right_idx_364 : tensor<256x2x1xi32> -> tensor<32x16xi32> loc(#loc186)
|
| 586 |
+
%cond_367 = arith.cmpi slt, %ileft_354, %iright_355 : tensor<32x16xi32> loc(#loc187)
|
| 587 |
+
%eq_368 = arith.cmpi eq, %ileft_354, %iright_355 : tensor<32x16xi32> loc(#loc188)
|
| 588 |
+
%cond_369 = arith.cmpi sgt, %left_idx_365, %right_idx_366 : tensor<32x16xi32> loc(#loc189)
|
| 589 |
+
%cond_370 = arith.andi %eq_368, %cond_369 : tensor<32x16xi1> loc(#loc190)
|
| 590 |
+
%cond_371 = arith.ori %cond_367, %cond_370 : tensor<32x16xi1> loc(#loc191)
|
| 591 |
+
%new_idxs_372 = arith.xori %left_idx_365, %right_idx_366 : tensor<32x16xi32> loc(#loc197)
|
| 592 |
+
%new_idxs_373 = arith.select %cond_371, %new_idxs_372, %cst_0 : tensor<32x16xi1>, tensor<32x16xi32> loc(#loc198)
|
| 593 |
+
%new_idxs_374 = arith.xori %new_idxs_344, %new_idxs_373 : tensor<32x16xi32> loc(#loc199)
|
| 594 |
+
%tmp7 = arith.extsi %tmp0_19 : tensor<32x16xi32> to tensor<32x16xi64> loc(#loc149)
|
| 595 |
+
%tmp10_375 = arith.select %tmp0_18, %tmp7, %tmp10 : tensor<32x16xi1>, tensor<32x16xi64> loc(#loc86)
|
| 596 |
+
%tmp11 = "tt.reduce"(%tmp10_375) <{axis = 1 : i32}> ({
|
| 597 |
+
^bb0(%tmp11_377: i64 loc(callsite(#loc1 at #loc150)), %tmp11_378: i64 loc(callsite(#loc1 at #loc150))):
|
| 598 |
+
%tmp11_379 = arith.addi %tmp11_377, %tmp11_378 : i64 loc(#loc200)
|
| 599 |
+
tt.reduce.return %tmp11_379 : i64 loc(#loc160)
|
| 600 |
+
}) : (tensor<32x16xi64>) -> tensor<32xi64> loc(#loc160)
|
| 601 |
+
%tmp11_376 = tt.expand_dims %tmp11 {axis = 1 : i32} : tensor<32xi64> -> tensor<32x1xi64> loc(#loc151)
|
| 602 |
+
%tmp14 = arith.trunci %tmp11_376 : tensor<32x1xi64> to tensor<32x1xi32> loc(#loc152)
|
| 603 |
+
%0 = arith.muli %xindex_6, %cst_2 : tensor<32x1xi32> loc(#loc73)
|
| 604 |
+
%1 = tt.broadcast %r0_index_8 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc74)
|
| 605 |
+
%2 = tt.broadcast %0 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc74)
|
| 606 |
+
%3 = arith.addi %1, %2 : tensor<32x16xi32> loc(#loc74)
|
| 607 |
+
%4 = tt.splat %out_ptr2 : !tt.ptr<i32> -> tensor<32x16x!tt.ptr<i32>> loc(#loc75)
|
| 608 |
+
%5 = tt.addptr %4, %3 : tensor<32x16x!tt.ptr<i32>>, tensor<32x16xi32> loc(#loc75)
|
| 609 |
+
tt.store %5, %new_idxs_374, %tmp0_18 : tensor<32x16x!tt.ptr<i32>> loc(#loc76)
|
| 610 |
+
%6 = tt.splat %out_ptr3 : !tt.ptr<i32> -> tensor<32x1x!tt.ptr<i32>> loc(#loc77)
|
| 611 |
+
%7 = tt.addptr %6, %xindex_6 : tensor<32x1x!tt.ptr<i32>>, tensor<32x1xi32> loc(#loc77)
|
| 612 |
+
tt.store %7, %tmp14, %xmask_7 : tensor<32x1x!tt.ptr<i32>> loc(#loc78)
|
| 613 |
+
tt.return loc(#loc79)
|
| 614 |
+
} loc(#loc)
|
| 615 |
+
} loc(#loc)
|
| 616 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":44:34)
|
| 617 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":36:49)
|
| 618 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":36:38)
|
| 619 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":26:21)
|
| 620 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":24:28)
|
| 621 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":24:33)
|
| 622 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":25:36)
|
| 623 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":25:44)
|
| 624 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":25:23)
|
| 625 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":27:28)
|
| 626 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":27:38)
|
| 627 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":33:19)
|
| 628 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":34:19)
|
| 629 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":36:35)
|
| 630 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":36:45)
|
| 631 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":36:30)
|
| 632 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":36:54)
|
| 633 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":38:19)
|
| 634 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":40:33)
|
| 635 |
+
#loc22 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:41)
|
| 636 |
+
#loc24 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:44)
|
| 637 |
+
#loc25 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:60)
|
| 638 |
+
#loc26 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":627:68)
|
| 639 |
+
#loc27 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":533:22)
|
| 640 |
+
#loc29 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":537:21)
|
| 641 |
+
#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:40)
|
| 642 |
+
#loc31 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 643 |
+
#loc33 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 644 |
+
#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:65)
|
| 645 |
+
#loc35 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":538:78)
|
| 646 |
+
#loc36 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:41)
|
| 647 |
+
#loc38 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:67)
|
| 648 |
+
#loc39 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":539:80)
|
| 649 |
+
#loc40 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":540:30)
|
| 650 |
+
#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":541:32)
|
| 651 |
+
#loc42 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":546:29)
|
| 652 |
+
#loc43 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:36)
|
| 653 |
+
#loc44 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:23)
|
| 654 |
+
#loc45 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":290:25)
|
| 655 |
+
#loc47 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:53)
|
| 656 |
+
#loc48 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":548:66)
|
| 657 |
+
#loc49 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:37)
|
| 658 |
+
#loc50 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:23)
|
| 659 |
+
#loc52 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:54)
|
| 660 |
+
#loc53 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":551:67)
|
| 661 |
+
#loc54 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":553:36)
|
| 662 |
+
#loc55 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":554:38)
|
| 663 |
+
#loc56 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":574:22)
|
| 664 |
+
#loc57 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":591:21)
|
| 665 |
+
#loc58 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:40)
|
| 666 |
+
#loc59 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:29)
|
| 667 |
+
#loc60 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":594:23)
|
| 668 |
+
#loc61 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":599:19)
|
| 669 |
+
#loc62 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":599:28)
|
| 670 |
+
#loc63 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:38)
|
| 671 |
+
#loc64 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:46)
|
| 672 |
+
#loc65 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":600:15)
|
| 673 |
+
#loc66 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:48)
|
| 674 |
+
#loc67 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:59)
|
| 675 |
+
#loc68 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":601:22)
|
| 676 |
+
#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":42:19)
|
| 677 |
+
#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":45:29)
|
| 678 |
+
#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":48:21)
|
| 679 |
+
#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":49:35)
|
| 680 |
+
#loc74 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":49:32)
|
| 681 |
+
#loc75 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":49:25)
|
| 682 |
+
#loc76 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":49:47)
|
| 683 |
+
#loc77 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":50:25)
|
| 684 |
+
#loc78 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":50:37)
|
| 685 |
+
#loc79 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/nj/cnjktwj7h4iwx4zghbum5atne46yt4ce4t5jnkkvyag35pn7glnh.py":50:4)
|
| 686 |
+
#loc85 = loc(callsite(#loc1 at #loc2))
|
| 687 |
+
#loc86 = loc("tmp10"(#loc3))
|
| 688 |
+
#loc87 = loc("tmp0"(#loc4))
|
| 689 |
+
#loc88 = loc("tmp0"(#loc5))
|
| 690 |
+
#loc89 = loc("xmask"(#loc6))
|
| 691 |
+
#loc90 = loc("xoffset"(#loc7))
|
| 692 |
+
#loc91 = loc("xoffset"(#loc8))
|
| 693 |
+
#loc92 = loc("xindex"(#loc9))
|
| 694 |
+
#loc93 = loc("xindex"(#loc10))
|
| 695 |
+
#loc94 = loc("xindex"(#loc11))
|
| 696 |
+
#loc95 = loc("r0_index"(#loc12))
|
| 697 |
+
#loc96 = loc("r0_index"(#loc13))
|
| 698 |
+
#loc97 = loc("x0"(#loc14))
|
| 699 |
+
#loc98 = loc("x1"(#loc15))
|
| 700 |
+
#loc99 = loc("tmp0"(#loc16))
|
| 701 |
+
#loc100 = loc("tmp0"(#loc17))
|
| 702 |
+
#loc101 = loc("tmp0"(#loc18))
|
| 703 |
+
#loc102 = loc("tmp0"(#loc19))
|
| 704 |
+
#loc103 = loc("tmp2"(#loc20))
|
| 705 |
+
#loc104 = loc("tmp4"(#loc21))
|
| 706 |
+
#loc105 = loc("flip"(#loc22))
|
| 707 |
+
#loc107 = loc("flip"(#loc24))
|
| 708 |
+
#loc108 = loc("flip"(#loc25))
|
| 709 |
+
#loc109 = loc("flip"(#loc26))
|
| 710 |
+
#loc110 = loc("y"(#loc27))
|
| 711 |
+
#loc111 = loc("left_mask"(#loc29))
|
| 712 |
+
#loc112 = loc("ileft"(#loc30))
|
| 713 |
+
#loc114 = loc("ileft"(#loc34))
|
| 714 |
+
#loc115 = loc("ileft"(#loc35))
|
| 715 |
+
#loc116 = loc("iright"(#loc36))
|
| 716 |
+
#loc118 = loc("iright"(#loc38))
|
| 717 |
+
#loc119 = loc("iright"(#loc39))
|
| 718 |
+
#loc120 = loc("ileft"(#loc40))
|
| 719 |
+
#loc121 = loc("iright"(#loc41))
|
| 720 |
+
#loc122 = loc("y_idx"(#loc42))
|
| 721 |
+
#loc123 = loc("left_idx"(#loc43))
|
| 722 |
+
#loc124 = loc("left_idx"(#loc44))
|
| 723 |
+
#loc125 = loc("input"(#loc45))
|
| 724 |
+
#loc127 = loc("left_idx"(#loc47))
|
| 725 |
+
#loc128 = loc("left_idx"(#loc48))
|
| 726 |
+
#loc129 = loc("right_idx"(#loc49))
|
| 727 |
+
#loc130 = loc("right_idx"(#loc50))
|
| 728 |
+
#loc132 = loc("right_idx"(#loc52))
|
| 729 |
+
#loc133 = loc("right_idx"(#loc53))
|
| 730 |
+
#loc134 = loc("left_idx"(#loc54))
|
| 731 |
+
#loc135 = loc("right_idx"(#loc55))
|
| 732 |
+
#loc136 = loc("cond"(#loc56))
|
| 733 |
+
#loc137 = loc("eq"(#loc57))
|
| 734 |
+
#loc138 = loc("cond"(#loc58))
|
| 735 |
+
#loc139 = loc("cond"(#loc59))
|
| 736 |
+
#loc140 = loc("cond"(#loc60))
|
| 737 |
+
#loc141 = loc("cond"(#loc61))
|
| 738 |
+
#loc142 = loc("cond"(#loc62))
|
| 739 |
+
#loc143 = loc("ret"(#loc63))
|
| 740 |
+
#loc144 = loc("ret"(#loc64))
|
| 741 |
+
#loc145 = loc("ret"(#loc65))
|
| 742 |
+
#loc146 = loc("new_idxs"(#loc66))
|
| 743 |
+
#loc147 = loc("new_idxs"(#loc67))
|
| 744 |
+
#loc148 = loc("new_idxs"(#loc68))
|
| 745 |
+
#loc149 = loc("tmp7"(#loc69))
|
| 746 |
+
#loc151 = loc("tmp11"(#loc71))
|
| 747 |
+
#loc152 = loc("tmp14"(#loc72))
|
| 748 |
+
#loc153 = loc(callsite(#loc105 at #loc106))
|
| 749 |
+
#loc154 = loc(callsite(#loc107 at #loc106))
|
| 750 |
+
#loc155 = loc(callsite(#loc108 at #loc106))
|
| 751 |
+
#loc156 = loc(callsite(#loc109 at #loc106))
|
| 752 |
+
#loc158 = loc("cond"(#loc136))
|
| 753 |
+
#loc159 = loc("eq"(#loc137))
|
| 754 |
+
#loc160 = loc(callsite(#loc31 at #loc150))
|
| 755 |
+
#loc162 = loc(callsite(#loc110 at #loc157))
|
| 756 |
+
#loc163 = loc(callsite(#loc111 at #loc157))
|
| 757 |
+
#loc164 = loc(callsite(#loc112 at #loc157))
|
| 758 |
+
#loc166 = loc(callsite(#loc114 at #loc157))
|
| 759 |
+
#loc167 = loc(callsite(#loc115 at #loc157))
|
| 760 |
+
#loc168 = loc(callsite(#loc116 at #loc157))
|
| 761 |
+
#loc170 = loc(callsite(#loc118 at #loc157))
|
| 762 |
+
#loc171 = loc(callsite(#loc119 at #loc157))
|
| 763 |
+
#loc172 = loc(callsite(#loc120 at #loc157))
|
| 764 |
+
#loc173 = loc(callsite(#loc121 at #loc157))
|
| 765 |
+
#loc174 = loc(callsite(#loc122 at #loc157))
|
| 766 |
+
#loc175 = loc(callsite(#loc123 at #loc157))
|
| 767 |
+
#loc176 = loc(callsite(#loc124 at #loc157))
|
| 768 |
+
#loc178 = loc(callsite(#loc127 at #loc157))
|
| 769 |
+
#loc179 = loc(callsite(#loc128 at #loc157))
|
| 770 |
+
#loc180 = loc(callsite(#loc129 at #loc157))
|
| 771 |
+
#loc181 = loc(callsite(#loc130 at #loc157))
|
| 772 |
+
#loc183 = loc(callsite(#loc132 at #loc157))
|
| 773 |
+
#loc184 = loc(callsite(#loc133 at #loc157))
|
| 774 |
+
#loc185 = loc(callsite(#loc134 at #loc157))
|
| 775 |
+
#loc186 = loc(callsite(#loc135 at #loc157))
|
| 776 |
+
#loc187 = loc(callsite(#loc158 at #loc157))
|
| 777 |
+
#loc188 = loc(callsite(#loc159 at #loc157))
|
| 778 |
+
#loc189 = loc(callsite(#loc138 at #loc157))
|
| 779 |
+
#loc190 = loc(callsite(#loc139 at #loc157))
|
| 780 |
+
#loc191 = loc(callsite(#loc140 at #loc157))
|
| 781 |
+
#loc192 = loc(callsite(#loc141 at #loc157))
|
| 782 |
+
#loc193 = loc(callsite(#loc142 at #loc157))
|
| 783 |
+
#loc194 = loc(callsite(#loc143 at #loc157))
|
| 784 |
+
#loc195 = loc(callsite(#loc144 at #loc157))
|
| 785 |
+
#loc196 = loc(callsite(#loc145 at #loc157))
|
| 786 |
+
#loc197 = loc(callsite(#loc146 at #loc157))
|
| 787 |
+
#loc198 = loc(callsite(#loc147 at #loc157))
|
| 788 |
+
#loc199 = loc(callsite(#loc148 at #loc157))
|
| 789 |
+
#loc200 = loc(callsite(#loc33 at #loc160))
|
| 790 |
+
#loc201 = loc(callsite(#loc31 at #loc165))
|
| 791 |
+
#loc203 = loc(callsite(#loc31 at #loc169))
|
| 792 |
+
#loc205 = loc(callsite(#loc125 at #loc177))
|
| 793 |
+
#loc206 = loc(callsite(#loc31 at #loc177))
|
| 794 |
+
#loc208 = loc(callsite(#loc125 at #loc182))
|
| 795 |
+
#loc209 = loc(callsite(#loc31 at #loc182))
|
| 796 |
+
#loc211 = loc(callsite(#loc33 at #loc201))
|
| 797 |
+
#loc212 = loc(callsite(#loc33 at #loc203))
|
| 798 |
+
#loc213 = loc(callsite(#loc33 at #loc206))
|
| 799 |
+
#loc214 = loc(callsite(#loc33 at #loc209))
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/__grp__triton_poi_fused_new_zeros_1.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_poi_fused_new_zeros_1.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.source", "triton_poi_fused_new_zeros_1.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ttir", "triton_poi_fused_new_zeros_1.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ttgir", "triton_poi_fused_new_zeros_1.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.llir", "triton_poi_fused_new_zeros_1.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ptx", "triton_poi_fused_new_zeros_1.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.cubin", "triton_poi_fused_new_zeros_1.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.cubin
ADDED
|
Binary file (5.38 kB). View file
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"hash": "e52d86863b53ed033419ead08893f8066d383f827ed32f19d25db032f55fb4c0", "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_new_zeros_1"}
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.llir
ADDED
|
@@ -0,0 +1,47 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
; ModuleID = 'LLVMDialectModule'
|
| 2 |
+
source_filename = "LLVMDialectModule"
|
| 3 |
+
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"
|
| 4 |
+
|
| 5 |
+
; Function Attrs: nounwind
|
| 6 |
+
define ptx_kernel void @triton_poi_fused_new_zeros_1(ptr addrspace(1) %0, i32 %1, ptr addrspace(1) readnone captures(none) %2, ptr addrspace(1) readnone captures(none) %3) local_unnamed_addr #0 !dbg !4 {
|
| 7 |
+
%5 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !7
|
| 8 |
+
%6 = shl i32 %5, 8, !dbg !8
|
| 9 |
+
%7 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !9
|
| 10 |
+
%8 = shl nuw nsw i32 %7, 1, !dbg !9
|
| 11 |
+
%9 = and i32 %8, 254, !dbg !9
|
| 12 |
+
%10 = or disjoint i32 %9, %6, !dbg !10
|
| 13 |
+
%11 = icmp slt i32 %10, 2176, !dbg !11
|
| 14 |
+
%12 = sext i32 %10 to i64, !dbg !12
|
| 15 |
+
%13 = getelementptr i32, ptr addrspace(1) %0, i64 %12, !dbg !12
|
| 16 |
+
tail call void asm sideeffect "@$3 st.global.v2.b32 [ $2 + 0 ], { $0, $1 };", "r,r,l,b"(i32 0, i32 0, ptr addrspace(1) %13, i1 %11) #2, !dbg !13
|
| 17 |
+
ret void, !dbg !14
|
| 18 |
+
}
|
| 19 |
+
|
| 20 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 21 |
+
declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1
|
| 22 |
+
|
| 23 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 24 |
+
declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
|
| 25 |
+
|
| 26 |
+
attributes #0 = { nounwind "nvvm.reqntid"="128" }
|
| 27 |
+
attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
|
| 28 |
+
attributes #2 = { nounwind }
|
| 29 |
+
|
| 30 |
+
!llvm.dbg.cu = !{!0}
|
| 31 |
+
!llvm.module.flags = !{!2, !3}
|
| 32 |
+
|
| 33 |
+
!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
|
| 34 |
+
!1 = !DIFile(filename: "c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f")
|
| 35 |
+
!2 = !{i32 2, !"Debug Info Version", i32 3}
|
| 36 |
+
!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
|
| 37 |
+
!4 = distinct !DISubprogram(name: "triton_poi_fused_new_zeros_1", linkageName: "triton_poi_fused_new_zeros_1", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
|
| 38 |
+
!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
|
| 39 |
+
!6 = !{}
|
| 40 |
+
!7 = !DILocation(line: 20, column: 28, scope: !4)
|
| 41 |
+
!8 = !DILocation(line: 20, column: 33, scope: !4)
|
| 42 |
+
!9 = !DILocation(line: 21, column: 36, scope: !4)
|
| 43 |
+
!10 = !DILocation(line: 21, column: 23, scope: !4)
|
| 44 |
+
!11 = !DILocation(line: 22, column: 21, scope: !4)
|
| 45 |
+
!12 = !DILocation(line: 25, column: 25, scope: !4)
|
| 46 |
+
!13 = !DILocation(line: 25, column: 36, scope: !4)
|
| 47 |
+
!14 = !DILocation(line: 25, column: 4, scope: !4)
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ptx
ADDED
|
@@ -0,0 +1,207 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
//
|
| 2 |
+
// Generated by LLVM NVPTX Back-End
|
| 3 |
+
//
|
| 4 |
+
|
| 5 |
+
.version 8.7
|
| 6 |
+
.target sm_90a
|
| 7 |
+
.address_size 64
|
| 8 |
+
|
| 9 |
+
// .globl triton_poi_fused_new_zeros_1 // -- Begin function triton_poi_fused_new_zeros_1
|
| 10 |
+
// @triton_poi_fused_new_zeros_1
|
| 11 |
+
.visible .entry triton_poi_fused_new_zeros_1(
|
| 12 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_new_zeros_1_param_0,
|
| 13 |
+
.param .u32 triton_poi_fused_new_zeros_1_param_1,
|
| 14 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_new_zeros_1_param_2,
|
| 15 |
+
.param .u64 .ptr .global .align 1 triton_poi_fused_new_zeros_1_param_3
|
| 16 |
+
)
|
| 17 |
+
.reqntid 128
|
| 18 |
+
{
|
| 19 |
+
.reg .pred %p<2>;
|
| 20 |
+
.reg .b32 %r<9>;
|
| 21 |
+
.reg .b64 %rd<3>;
|
| 22 |
+
.loc 1 18 0 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:18:0
|
| 23 |
+
$L__func_begin0:
|
| 24 |
+
.loc 1 18 0 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:18:0
|
| 25 |
+
|
| 26 |
+
// %bb.0:
|
| 27 |
+
ld.param.b64 %rd2, [triton_poi_fused_new_zeros_1_param_0];
|
| 28 |
+
$L__tmp0:
|
| 29 |
+
.loc 1 20 28 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:20:28
|
| 30 |
+
mov.u32 %r3, %ctaid.x;
|
| 31 |
+
.loc 1 20 33 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:20:33
|
| 32 |
+
shl.b32 %r4, %r3, 8;
|
| 33 |
+
.loc 1 21 36 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:21:36
|
| 34 |
+
mov.u32 %r5, %tid.x;
|
| 35 |
+
shl.b32 %r6, %r5, 1;
|
| 36 |
+
and.b32 %r7, %r6, 254;
|
| 37 |
+
.loc 1 21 23 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:21:23
|
| 38 |
+
or.b32 %r8, %r7, %r4;
|
| 39 |
+
.loc 1 22 21 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:22:21
|
| 40 |
+
setp.lt.s32 %p1, %r8, 2176;
|
| 41 |
+
.loc 1 25 25 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:25:25
|
| 42 |
+
mad.wide.s32 %rd1, %r8, 4, %rd2;
|
| 43 |
+
mov.b32 %r1, 0;
|
| 44 |
+
.loc 1 25 36 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:25:36
|
| 45 |
+
// begin inline asm
|
| 46 |
+
@%p1 st.global.v2.b32 [ %rd1 + 0 ], { %r1, %r1 };
|
| 47 |
+
// end inline asm
|
| 48 |
+
.loc 1 25 4 // c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py:25:4
|
| 49 |
+
ret;
|
| 50 |
+
$L__tmp1:
|
| 51 |
+
$L__func_end0:
|
| 52 |
+
// -- End function
|
| 53 |
+
}
|
| 54 |
+
.file 1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py"
|
| 55 |
+
.section .debug_abbrev
|
| 56 |
+
{
|
| 57 |
+
.b8 1 // Abbreviation Code
|
| 58 |
+
.b8 17 // DW_TAG_compile_unit
|
| 59 |
+
.b8 0 // DW_CHILDREN_no
|
| 60 |
+
.b8 37 // DW_AT_producer
|
| 61 |
+
.b8 8 // DW_FORM_string
|
| 62 |
+
.b8 19 // DW_AT_language
|
| 63 |
+
.b8 5 // DW_FORM_data2
|
| 64 |
+
.b8 3 // DW_AT_name
|
| 65 |
+
.b8 8 // DW_FORM_string
|
| 66 |
+
.b8 16 // DW_AT_stmt_list
|
| 67 |
+
.b8 6 // DW_FORM_data4
|
| 68 |
+
.b8 27 // DW_AT_comp_dir
|
| 69 |
+
.b8 8 // DW_FORM_string
|
| 70 |
+
.b8 0 // EOM(1)
|
| 71 |
+
.b8 0 // EOM(2)
|
| 72 |
+
.b8 0 // EOM(3)
|
| 73 |
+
}
|
| 74 |
+
.section .debug_info
|
| 75 |
+
{
|
| 76 |
+
.b32 135 // Length of Unit
|
| 77 |
+
.b8 2 // DWARF version number
|
| 78 |
+
.b8 0
|
| 79 |
+
.b32 .debug_abbrev // Offset Into Abbrev. Section
|
| 80 |
+
.b8 8 // Address Size (in bytes)
|
| 81 |
+
.b8 1 // Abbrev [1] 0xb:0x80 DW_TAG_compile_unit
|
| 82 |
+
.b8 116 // DW_AT_producer
|
| 83 |
+
.b8 114
|
| 84 |
+
.b8 105
|
| 85 |
+
.b8 116
|
| 86 |
+
.b8 111
|
| 87 |
+
.b8 110
|
| 88 |
+
.b8 0
|
| 89 |
+
.b8 2 // DW_AT_language
|
| 90 |
+
.b8 0
|
| 91 |
+
.b8 99 // DW_AT_name
|
| 92 |
+
.b8 54
|
| 93 |
+
.b8 102
|
| 94 |
+
.b8 117
|
| 95 |
+
.b8 104
|
| 96 |
+
.b8 99
|
| 97 |
+
.b8 116
|
| 98 |
+
.b8 53
|
| 99 |
+
.b8 118
|
| 100 |
+
.b8 100
|
| 101 |
+
.b8 112
|
| 102 |
+
.b8 51
|
| 103 |
+
.b8 100
|
| 104 |
+
.b8 53
|
| 105 |
+
.b8 108
|
| 106 |
+
.b8 120
|
| 107 |
+
.b8 52
|
| 108 |
+
.b8 53
|
| 109 |
+
.b8 99
|
| 110 |
+
.b8 104
|
| 111 |
+
.b8 122
|
| 112 |
+
.b8 50
|
| 113 |
+
.b8 55
|
| 114 |
+
.b8 103
|
| 115 |
+
.b8 104
|
| 116 |
+
.b8 97
|
| 117 |
+
.b8 103
|
| 118 |
+
.b8 53
|
| 119 |
+
.b8 100
|
| 120 |
+
.b8 102
|
| 121 |
+
.b8 114
|
| 122 |
+
.b8 101
|
| 123 |
+
.b8 104
|
| 124 |
+
.b8 50
|
| 125 |
+
.b8 104
|
| 126 |
+
.b8 51
|
| 127 |
+
.b8 104
|
| 128 |
+
.b8 98
|
| 129 |
+
.b8 122
|
| 130 |
+
.b8 120
|
| 131 |
+
.b8 108
|
| 132 |
+
.b8 53
|
| 133 |
+
.b8 101
|
| 134 |
+
.b8 108
|
| 135 |
+
.b8 104
|
| 136 |
+
.b8 105
|
| 137 |
+
.b8 109
|
| 138 |
+
.b8 51
|
| 139 |
+
.b8 113
|
| 140 |
+
.b8 104
|
| 141 |
+
.b8 112
|
| 142 |
+
.b8 120
|
| 143 |
+
.b8 46
|
| 144 |
+
.b8 112
|
| 145 |
+
.b8 121
|
| 146 |
+
.b8 0
|
| 147 |
+
.b32 .debug_line // DW_AT_stmt_list
|
| 148 |
+
.b8 47 // DW_AT_comp_dir
|
| 149 |
+
.b8 119
|
| 150 |
+
.b8 111
|
| 151 |
+
.b8 114
|
| 152 |
+
.b8 107
|
| 153 |
+
.b8 115
|
| 154 |
+
.b8 112
|
| 155 |
+
.b8 97
|
| 156 |
+
.b8 99
|
| 157 |
+
.b8 101
|
| 158 |
+
.b8 47
|
| 159 |
+
.b8 104
|
| 160 |
+
.b8 97
|
| 161 |
+
.b8 110
|
| 162 |
+
.b8 114
|
| 163 |
+
.b8 117
|
| 164 |
+
.b8 105
|
| 165 |
+
.b8 47
|
| 166 |
+
.b8 83
|
| 167 |
+
.b8 112
|
| 168 |
+
.b8 101
|
| 169 |
+
.b8 99
|
| 170 |
+
.b8 70
|
| 171 |
+
.b8 111
|
| 172 |
+
.b8 114
|
| 173 |
+
.b8 103
|
| 174 |
+
.b8 101
|
| 175 |
+
.b8 45
|
| 176 |
+
.b8 101
|
| 177 |
+
.b8 120
|
| 178 |
+
.b8 116
|
| 179 |
+
.b8 47
|
| 180 |
+
.b8 99
|
| 181 |
+
.b8 97
|
| 182 |
+
.b8 99
|
| 183 |
+
.b8 104
|
| 184 |
+
.b8 101
|
| 185 |
+
.b8 47
|
| 186 |
+
.b8 99
|
| 187 |
+
.b8 111
|
| 188 |
+
.b8 109
|
| 189 |
+
.b8 112
|
| 190 |
+
.b8 105
|
| 191 |
+
.b8 108
|
| 192 |
+
.b8 101
|
| 193 |
+
.b8 100
|
| 194 |
+
.b8 95
|
| 195 |
+
.b8 107
|
| 196 |
+
.b8 101
|
| 197 |
+
.b8 114
|
| 198 |
+
.b8 110
|
| 199 |
+
.b8 101
|
| 200 |
+
.b8 108
|
| 201 |
+
.b8 115
|
| 202 |
+
.b8 47
|
| 203 |
+
.b8 54
|
| 204 |
+
.b8 102
|
| 205 |
+
.b8 0
|
| 206 |
+
}
|
| 207 |
+
.section .debug_macinfo { }
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.source
ADDED
|
@@ -0,0 +1,41 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":18:0)
|
| 2 |
+
#loc11 = loc("out_ptr0"(#loc))
|
| 3 |
+
#loc12 = loc("xnumel"(#loc))
|
| 4 |
+
module {
|
| 5 |
+
tt.func public @triton_poi_fused_new_zeros_1(%out_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc))) attributes {noinline = false} {
|
| 6 |
+
%xnumel_0 = arith.constant 2176 : i32 loc(#loc13)
|
| 7 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc14)
|
| 8 |
+
%xoffset_1 = arith.constant 256 : i32 loc(#loc15)
|
| 9 |
+
%xoffset_2 = arith.constant 256 : i32 loc(#loc15)
|
| 10 |
+
%xoffset_3 = arith.muli %xoffset, %xoffset_2 : i32 loc(#loc15)
|
| 11 |
+
%xindex = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32> loc(#loc16)
|
| 12 |
+
%xindex_4 = tt.splat %xoffset_3 : i32 -> tensor<256xi32> loc(#loc17)
|
| 13 |
+
%xindex_5 = arith.addi %xindex_4, %xindex : tensor<256xi32> loc(#loc17)
|
| 14 |
+
%xmask = arith.constant dense<2176> : tensor<256xi32> loc(#loc18)
|
| 15 |
+
%xmask_6 = arith.cmpi slt, %xindex_5, %xmask : tensor<256xi32> loc(#loc18)
|
| 16 |
+
%tmp0 = arith.constant 0 : i32 loc(#loc19)
|
| 17 |
+
%tmp0_7 = arith.constant dense<0> : tensor<1xi32> loc(#loc19)
|
| 18 |
+
%0 = tt.splat %out_ptr0 : !tt.ptr<i32> -> tensor<256x!tt.ptr<i32>> loc(#loc8)
|
| 19 |
+
%1 = tt.addptr %0, %xindex_5 : tensor<256x!tt.ptr<i32>>, tensor<256xi32> loc(#loc8)
|
| 20 |
+
%cst = arith.constant dense<0> : tensor<256xi32> loc(#loc9)
|
| 21 |
+
tt.store %1, %cst, %xmask_6 : tensor<256x!tt.ptr<i32>> loc(#loc9)
|
| 22 |
+
tt.return loc(#loc10)
|
| 23 |
+
} loc(#loc)
|
| 24 |
+
} loc(#loc)
|
| 25 |
+
#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":19:13)
|
| 26 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":20:28)
|
| 27 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":20:33)
|
| 28 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":21:36)
|
| 29 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":21:23)
|
| 30 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":22:21)
|
| 31 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":24:27)
|
| 32 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:25)
|
| 33 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:36)
|
| 34 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:4)
|
| 35 |
+
#loc13 = loc("xnumel"(#loc1))
|
| 36 |
+
#loc14 = loc("xoffset"(#loc2))
|
| 37 |
+
#loc15 = loc("xoffset"(#loc3))
|
| 38 |
+
#loc16 = loc("xindex"(#loc4))
|
| 39 |
+
#loc17 = loc("xindex"(#loc5))
|
| 40 |
+
#loc18 = loc("xmask"(#loc6))
|
| 41 |
+
#loc19 = loc("tmp0"(#loc7))
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ttgir
ADDED
|
@@ -0,0 +1,35 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#blocked = #ttg.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
|
| 2 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":18:0)
|
| 3 |
+
#loc10 = loc("out_ptr0"(#loc))
|
| 4 |
+
#loc11 = loc("xnumel"(#loc))
|
| 5 |
+
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
|
| 6 |
+
tt.func public @triton_poi_fused_new_zeros_1(%out_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc))) attributes {noinline = false} {
|
| 7 |
+
%cst = arith.constant dense<2176> : tensor<256xi32, #blocked> loc(#loc1)
|
| 8 |
+
%cst_0 = arith.constant dense<0> : tensor<256xi32, #blocked> loc(#loc1)
|
| 9 |
+
%c256_i32 = arith.constant 256 : i32 loc(#loc1)
|
| 10 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc12)
|
| 11 |
+
%xoffset_1 = arith.muli %xoffset, %c256_i32 : i32 loc(#loc13)
|
| 12 |
+
%xindex = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #blocked> loc(#loc14)
|
| 13 |
+
%xindex_2 = tt.splat %xoffset_1 : i32 -> tensor<256xi32, #blocked> loc(#loc15)
|
| 14 |
+
%xindex_3 = arith.addi %xindex_2, %xindex : tensor<256xi32, #blocked> loc(#loc15)
|
| 15 |
+
%xmask = arith.cmpi slt, %xindex_3, %cst : tensor<256xi32, #blocked> loc(#loc16)
|
| 16 |
+
%0 = tt.splat %out_ptr0 : !tt.ptr<i32> -> tensor<256x!tt.ptr<i32>, #blocked> loc(#loc7)
|
| 17 |
+
%1 = tt.addptr %0, %xindex_3 : tensor<256x!tt.ptr<i32>, #blocked>, tensor<256xi32, #blocked> loc(#loc7)
|
| 18 |
+
tt.store %1, %cst_0, %xmask : tensor<256x!tt.ptr<i32>, #blocked> loc(#loc8)
|
| 19 |
+
tt.return loc(#loc9)
|
| 20 |
+
} loc(#loc)
|
| 21 |
+
} loc(#loc)
|
| 22 |
+
#loc1 = loc(unknown)
|
| 23 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":20:28)
|
| 24 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":20:33)
|
| 25 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":21:36)
|
| 26 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":21:23)
|
| 27 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":22:21)
|
| 28 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:25)
|
| 29 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:36)
|
| 30 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:4)
|
| 31 |
+
#loc12 = loc("xoffset"(#loc2))
|
| 32 |
+
#loc13 = loc("xoffset"(#loc3))
|
| 33 |
+
#loc14 = loc("xindex"(#loc4))
|
| 34 |
+
#loc15 = loc("xindex"(#loc5))
|
| 35 |
+
#loc16 = loc("xmask"(#loc6))
|
SpecForge-ext/cache/compiled_kernels/triton/1/4UWYNBR3KPWQGNAZ5LIIRE7YAZWTQP4CP3JS6GOSLWYDF5K7WTAA/triton_poi_fused_new_zeros_1.ttir
ADDED
|
@@ -0,0 +1,34 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":18:0)
|
| 2 |
+
#loc10 = loc("out_ptr0"(#loc))
|
| 3 |
+
#loc11 = loc("xnumel"(#loc))
|
| 4 |
+
module {
|
| 5 |
+
tt.func public @triton_poi_fused_new_zeros_1(%out_ptr0: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr0"(#loc)), %xnumel: i32 {tt.divisibility = 16 : i32} loc("xnumel"(#loc))) attributes {noinline = false} {
|
| 6 |
+
%cst = arith.constant dense<0> : tensor<256xi32> loc(#loc1)
|
| 7 |
+
%xmask = arith.constant dense<2176> : tensor<256xi32> loc(#loc12)
|
| 8 |
+
%c256_i32 = arith.constant 256 : i32 loc(#loc3)
|
| 9 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc13)
|
| 10 |
+
%xoffset_0 = arith.muli %xoffset, %c256_i32 : i32 loc(#loc14)
|
| 11 |
+
%xindex = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32> loc(#loc15)
|
| 12 |
+
%xindex_1 = tt.splat %xoffset_0 : i32 -> tensor<256xi32> loc(#loc16)
|
| 13 |
+
%xindex_2 = arith.addi %xindex_1, %xindex : tensor<256xi32> loc(#loc16)
|
| 14 |
+
%xmask_3 = arith.cmpi slt, %xindex_2, %xmask : tensor<256xi32> loc(#loc12)
|
| 15 |
+
%0 = tt.splat %out_ptr0 : !tt.ptr<i32> -> tensor<256x!tt.ptr<i32>> loc(#loc8)
|
| 16 |
+
%1 = tt.addptr %0, %xindex_2 : tensor<256x!tt.ptr<i32>>, tensor<256xi32> loc(#loc8)
|
| 17 |
+
tt.store %1, %cst, %xmask_3 : tensor<256x!tt.ptr<i32>> loc(#loc1)
|
| 18 |
+
tt.return loc(#loc9)
|
| 19 |
+
} loc(#loc)
|
| 20 |
+
} loc(#loc)
|
| 21 |
+
#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:36)
|
| 22 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":22:21)
|
| 23 |
+
#loc3 = loc(unknown)
|
| 24 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":20:28)
|
| 25 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":20:33)
|
| 26 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":21:36)
|
| 27 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":21:23)
|
| 28 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:25)
|
| 29 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/6f/c6fuhct5vdp3d5lx45chz27ghag5dfreh2h3hbzxl5elhim3qhpx.py":25:4)
|
| 30 |
+
#loc12 = loc("xmask"(#loc2))
|
| 31 |
+
#loc13 = loc("xoffset"(#loc4))
|
| 32 |
+
#loc14 = loc("xoffset"(#loc5))
|
| 33 |
+
#loc15 = loc("xindex"(#loc6))
|
| 34 |
+
#loc16 = loc("xindex"(#loc7))
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/__grp__triton_red_fused__to_copy_mul_sum_0.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_red_fused__to_copy_mul_sum_0.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.source", "triton_red_fused__to_copy_mul_sum_0.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ttir", "triton_red_fused__to_copy_mul_sum_0.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ttgir", "triton_red_fused__to_copy_mul_sum_0.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.llir", "triton_red_fused__to_copy_mul_sum_0.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ptx", "triton_red_fused__to_copy_mul_sum_0.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.cubin", "triton_red_fused__to_copy_mul_sum_0.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.cubin
ADDED
|
Binary file (27.6 kB). View file
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"hash": "eab00da2e4595f249937e6537c90d9a940b19c806c3f885fbde47bd85402f17a", "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": 512, "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_mul_sum_0"}
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.llir
ADDED
|
@@ -0,0 +1,256 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
; ModuleID = 'LLVMDialectModule'
|
| 2 |
+
source_filename = "LLVMDialectModule"
|
| 3 |
+
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"
|
| 4 |
+
|
| 5 |
+
@global_smem = external addrspace(3) global [0 x i8], align 16
|
| 6 |
+
|
| 7 |
+
; Function Attrs: nounwind
|
| 8 |
+
define ptx_kernel void @triton_red_fused__to_copy_mul_sum_0(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, i64 %4, i64 %5, i64 %6, i32 %7, i32 %8, ptr addrspace(1) readnone captures(none) %9, ptr addrspace(1) readnone captures(none) %10) local_unnamed_addr #0 !dbg !4 {
|
| 9 |
+
%12 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !7
|
| 10 |
+
%13 = shl i32 %12, 6, !dbg !8
|
| 11 |
+
%14 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !9
|
| 12 |
+
%15 = and i32 %14, 63, !dbg !9
|
| 13 |
+
%16 = or disjoint i32 %13, %15, !dbg !10
|
| 14 |
+
%17 = icmp slt i32 %16, %7, !dbg !11
|
| 15 |
+
%18 = lshr i32 %14, 6, !dbg !12
|
| 16 |
+
%.lobit = and i32 %18, 1, !dbg !12
|
| 17 |
+
%19 = sext i32 %16 to i64, !dbg !13
|
| 18 |
+
%.frozen = freeze i64 %4, !dbg !13
|
| 19 |
+
%20 = sdiv i64 %19, %.frozen, !dbg !13
|
| 20 |
+
%21 = mul i64 %20, %.frozen, !dbg !14
|
| 21 |
+
%.decomposed = sub i64 %19, %21, !dbg !14
|
| 22 |
+
%22 = mul i64 %6, %5, !dbg !15
|
| 23 |
+
%23 = add i64 %22, 31, !dbg !16
|
| 24 |
+
%24 = sdiv i64 %23, 32, !dbg !17
|
| 25 |
+
%25 = mul i64 %20, %24, !dbg !18
|
| 26 |
+
%26 = icmp sgt i32 %8, 0, !dbg !19
|
| 27 |
+
br i1 %26, label %.lr.ph.preheader, label %._crit_edge, !dbg !19
|
| 28 |
+
|
| 29 |
+
.lr.ph.preheader: ; preds = %11
|
| 30 |
+
%27 = insertelement <4 x i32> poison, i32 %8, i64 0
|
| 31 |
+
%28 = shufflevector <4 x i32> %27, <4 x i32> poison, <4 x i32> zeroinitializer
|
| 32 |
+
%29 = insertelement <4 x i64> poison, i64 %25, i64 0
|
| 33 |
+
%30 = shufflevector <4 x i64> %29, <4 x i64> poison, <4 x i32> zeroinitializer
|
| 34 |
+
%31 = insertelement <4 x i64> poison, i64 %22, i64 0
|
| 35 |
+
%32 = shufflevector <4 x i64> %31, <4 x i64> poison, <4 x i32> zeroinitializer
|
| 36 |
+
%33 = insertelement <4 x i1> poison, i1 %17, i64 0
|
| 37 |
+
%34 = shufflevector <4 x i1> %33, <4 x i1> poison, <4 x i32> zeroinitializer
|
| 38 |
+
br label %.lr.ph, !dbg !19
|
| 39 |
+
|
| 40 |
+
.lr.ph: ; preds = %.lr.ph.preheader, %.lr.ph
|
| 41 |
+
%35 = phi i32 [ %132, %.lr.ph ], [ 0, %.lr.ph.preheader ]
|
| 42 |
+
%36 = phi <4 x float> [ %131, %.lr.ph ], [ zeroinitializer, %.lr.ph.preheader ]
|
| 43 |
+
%37 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 44 |
+
%38 = or disjoint i32 %35, %.lobit, !dbg !21
|
| 45 |
+
%39 = or disjoint i32 %38, 2, !dbg !21
|
| 46 |
+
%40 = or disjoint i32 %38, 4, !dbg !21
|
| 47 |
+
%41 = or disjoint i32 %38, 6, !dbg !21
|
| 48 |
+
%42 = insertelement <4 x i32> poison, i32 %38, i64 0, !dbg !22
|
| 49 |
+
%43 = insertelement <4 x i32> %42, i32 %39, i64 1, !dbg !22
|
| 50 |
+
%44 = insertelement <4 x i32> %43, i32 %40, i64 2, !dbg !22
|
| 51 |
+
%45 = insertelement <4 x i32> %44, i32 %41, i64 3, !dbg !22
|
| 52 |
+
%46 = icmp slt <4 x i32> %45, %28, !dbg !22
|
| 53 |
+
%47 = sext <4 x i32> %45 to <4 x i64>, !dbg !23
|
| 54 |
+
%48 = add <4 x i64> %30, %47, !dbg !23
|
| 55 |
+
%49 = icmp slt <4 x i64> %48, %32, !dbg !24
|
| 56 |
+
%50 = extractelement <4 x i64> %48, i64 0, !dbg !25
|
| 57 |
+
%51 = srem i64 %50, %22, !dbg !25
|
| 58 |
+
%52 = extractelement <4 x i64> %48, i64 1, !dbg !25
|
| 59 |
+
%53 = srem i64 %52, %22, !dbg !25
|
| 60 |
+
%54 = extractelement <4 x i64> %48, i64 2, !dbg !25
|
| 61 |
+
%55 = srem i64 %54, %22, !dbg !25
|
| 62 |
+
%56 = extractelement <4 x i64> %48, i64 3, !dbg !25
|
| 63 |
+
%57 = srem i64 %56, %22, !dbg !25
|
| 64 |
+
%58 = mul i64 %51, %4, !dbg !26
|
| 65 |
+
%59 = mul i64 %53, %4, !dbg !26
|
| 66 |
+
%60 = mul i64 %55, %4, !dbg !26
|
| 67 |
+
%61 = mul i64 %57, %4, !dbg !26
|
| 68 |
+
%62 = add i64 %58, %.decomposed, !dbg !27
|
| 69 |
+
%63 = add i64 %59, %.decomposed, !dbg !27
|
| 70 |
+
%64 = add i64 %60, %.decomposed, !dbg !27
|
| 71 |
+
%65 = add i64 %61, %.decomposed, !dbg !27
|
| 72 |
+
%66 = getelementptr bfloat, ptr addrspace(1) %0, i64 %62, !dbg !28
|
| 73 |
+
%67 = getelementptr bfloat, ptr addrspace(1) %0, i64 %63, !dbg !28
|
| 74 |
+
%68 = getelementptr bfloat, ptr addrspace(1) %0, i64 %64, !dbg !28
|
| 75 |
+
%69 = getelementptr bfloat, ptr addrspace(1) %0, i64 %65, !dbg !28
|
| 76 |
+
%foldExtExtBinop = and <4 x i1> %46, %49, !dbg !29
|
| 77 |
+
%70 = extractelement <4 x i1> %foldExtExtBinop, i64 0, !dbg !29
|
| 78 |
+
%foldExtExtBinop8 = and <4 x i1> %46, %49, !dbg !29
|
| 79 |
+
%71 = extractelement <4 x i1> %foldExtExtBinop8, i64 1, !dbg !29
|
| 80 |
+
%foldExtExtBinop10 = and <4 x i1> %46, %49, !dbg !29
|
| 81 |
+
%72 = extractelement <4 x i1> %foldExtExtBinop10, i64 2, !dbg !29
|
| 82 |
+
%foldExtExtBinop12 = and <4 x i1> %46, %49, !dbg !29
|
| 83 |
+
%73 = extractelement <4 x i1> %foldExtExtBinop12, i64 3, !dbg !29
|
| 84 |
+
%74 = and i1 %17, %70, !dbg !30
|
| 85 |
+
%75 = and i1 %17, %71, !dbg !30
|
| 86 |
+
%76 = and i1 %17, %72, !dbg !30
|
| 87 |
+
%77 = and i1 %17, %73, !dbg !30
|
| 88 |
+
%78 = 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) %66, i64 %37, i1 %74) #4, !dbg !20
|
| 89 |
+
%79 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 90 |
+
%80 = 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) %67, i64 %79, i1 %75) #4, !dbg !20
|
| 91 |
+
%81 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 92 |
+
%82 = 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) %68, i64 %81, i1 %76) #4, !dbg !20
|
| 93 |
+
%83 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !20
|
| 94 |
+
%84 = 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) %69, i64 %83, i1 %77) #4, !dbg !20
|
| 95 |
+
%85 = insertelement <4 x i16> poison, i16 %78, i64 0, !dbg !20
|
| 96 |
+
%86 = insertelement <4 x i16> %85, i16 %80, i64 1, !dbg !20
|
| 97 |
+
%87 = insertelement <4 x i16> %86, i16 %82, i64 2, !dbg !20
|
| 98 |
+
%88 = insertelement <4 x i16> %87, i16 %84, i64 3, !dbg !20
|
| 99 |
+
%89 = bitcast <4 x i16> %88 to <4 x bfloat>, !dbg !20
|
| 100 |
+
%90 = fpext <4 x bfloat> %89 to <4 x float>, !dbg !31
|
| 101 |
+
%91 = getelementptr bfloat, ptr addrspace(1) %1, i64 %62, !dbg !32
|
| 102 |
+
%92 = getelementptr bfloat, ptr addrspace(1) %1, i64 %63, !dbg !32
|
| 103 |
+
%93 = getelementptr bfloat, ptr addrspace(1) %1, i64 %64, !dbg !32
|
| 104 |
+
%94 = getelementptr bfloat, ptr addrspace(1) %1, i64 %65, !dbg !32
|
| 105 |
+
%95 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !33
|
| 106 |
+
%96 = 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) %91, i64 %95, i1 %74) #4, !dbg !33
|
| 107 |
+
%97 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !33
|
| 108 |
+
%98 = 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) %92, i64 %97, i1 %75) #4, !dbg !33
|
| 109 |
+
%99 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !33
|
| 110 |
+
%100 = 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) %93, i64 %99, i1 %76) #4, !dbg !33
|
| 111 |
+
%101 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !33
|
| 112 |
+
%102 = 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) %94, i64 %101, i1 %77) #4, !dbg !33
|
| 113 |
+
%103 = insertelement <4 x i16> poison, i16 %96, i64 0, !dbg !33
|
| 114 |
+
%104 = insertelement <4 x i16> %103, i16 %98, i64 1, !dbg !33
|
| 115 |
+
%105 = insertelement <4 x i16> %104, i16 %100, i64 2, !dbg !33
|
| 116 |
+
%106 = insertelement <4 x i16> %105, i16 %102, i64 3, !dbg !33
|
| 117 |
+
%107 = bitcast <4 x i16> %106 to <4 x bfloat>, !dbg !33
|
| 118 |
+
%108 = fpext <4 x bfloat> %107 to <4 x float>, !dbg !34
|
| 119 |
+
%109 = getelementptr float, ptr addrspace(1) %2, i64 %51, !dbg !35
|
| 120 |
+
%110 = getelementptr float, ptr addrspace(1) %2, i64 %53, !dbg !35
|
| 121 |
+
%111 = getelementptr float, ptr addrspace(1) %2, i64 %55, !dbg !35
|
| 122 |
+
%112 = getelementptr float, ptr addrspace(1) %2, i64 %57, !dbg !35
|
| 123 |
+
%113 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !36
|
| 124 |
+
%114 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b32 { $0 }, [ $2 + 0 ], $3;", "=r,r,l,l,b"(i32 0, ptr addrspace(1) %109, i64 %113, i1 %74) #4, !dbg !36
|
| 125 |
+
%115 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !36
|
| 126 |
+
%116 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b32 { $0 }, [ $2 + 0 ], $3;", "=r,r,l,l,b"(i32 0, ptr addrspace(1) %110, i64 %115, i1 %75) #4, !dbg !36
|
| 127 |
+
%117 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !36
|
| 128 |
+
%118 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b32 { $0 }, [ $2 + 0 ], $3;", "=r,r,l,l,b"(i32 0, ptr addrspace(1) %111, i64 %117, i1 %76) #4, !dbg !36
|
| 129 |
+
%119 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", "=l"() #4, !dbg !36
|
| 130 |
+
%120 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$4 ld.global.L1::evict_last.L2::cache_hint.b32 { $0 }, [ $2 + 0 ], $3;", "=r,r,l,l,b"(i32 0, ptr addrspace(1) %112, i64 %119, i1 %77) #4, !dbg !36
|
| 131 |
+
%121 = insertelement <4 x i32> poison, i32 %114, i64 0, !dbg !36
|
| 132 |
+
%122 = insertelement <4 x i32> %121, i32 %116, i64 1, !dbg !36
|
| 133 |
+
%123 = insertelement <4 x i32> %122, i32 %118, i64 2, !dbg !36
|
| 134 |
+
%124 = insertelement <4 x i32> %123, i32 %120, i64 3, !dbg !36
|
| 135 |
+
%125 = bitcast <4 x i32> %124 to <4 x float>, !dbg !36
|
| 136 |
+
%126 = fmul <4 x float> %108, %125, !dbg !37
|
| 137 |
+
%127 = fmul <4 x float> %126, %90, !dbg !38
|
| 138 |
+
%128 = select <4 x i1> %49, <4 x float> %127, <4 x float> zeroinitializer, !dbg !39
|
| 139 |
+
%129 = fadd <4 x float> %36, %128, !dbg !40
|
| 140 |
+
%130 = and <4 x i1> %34, %46, !dbg !41
|
| 141 |
+
%131 = select <4 x i1> %130, <4 x float> %129, <4 x float> %36, !dbg !42
|
| 142 |
+
%132 = add i32 %35, 8, !dbg !19
|
| 143 |
+
%133 = icmp slt i32 %132, %8, !dbg !19
|
| 144 |
+
br i1 %133, label %.lr.ph, label %._crit_edge.loopexit, !dbg !19
|
| 145 |
+
|
| 146 |
+
._crit_edge.loopexit: ; preds = %.lr.ph
|
| 147 |
+
%shift = shufflevector <4 x float> %131, <4 x float> poison, <4 x i32> <i32 1, i32 poison, i32 poison, i32 poison>, !dbg !43
|
| 148 |
+
%foldExtExtBinop14 = fadd <4 x float> %131, %shift, !dbg !43
|
| 149 |
+
%shift16 = shufflevector <4 x float> %131, <4 x float> poison, <4 x i32> <i32 2, i32 poison, i32 poison, i32 poison>, !dbg !43
|
| 150 |
+
%foldExtExtBinop17 = fadd <4 x float> %shift16, %foldExtExtBinop14, !dbg !43
|
| 151 |
+
%shift19 = shufflevector <4 x float> %131, <4 x float> poison, <4 x i32> <i32 3, i32 poison, i32 poison, i32 poison>, !dbg !43
|
| 152 |
+
%foldExtExtBinop20 = fadd <4 x float> %shift19, %foldExtExtBinop17, !dbg !43
|
| 153 |
+
%134 = extractelement <4 x float> %foldExtExtBinop20, i64 0, !dbg !43
|
| 154 |
+
%135 = bitcast float %134 to <1 x i32>, !dbg !47
|
| 155 |
+
br label %._crit_edge, !dbg !43
|
| 156 |
+
|
| 157 |
+
._crit_edge: ; preds = %._crit_edge.loopexit, %11
|
| 158 |
+
%136 = phi <1 x i32> [ zeroinitializer, %11 ], [ %135, %._crit_edge.loopexit ], !dbg !43
|
| 159 |
+
%.idx = shl nuw nsw i32 %15, 3, !dbg !47
|
| 160 |
+
%137 = getelementptr i8, ptr addrspace(3) @global_smem, i32 %.idx, !dbg !47
|
| 161 |
+
%138 = getelementptr float, ptr addrspace(3) %137, i32 %.lobit, !dbg !47
|
| 162 |
+
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %138, <1 x i32> %136, i1 true) #4, !dbg !47
|
| 163 |
+
tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !47
|
| 164 |
+
%139 = icmp samesign ult i32 %14, 128, !dbg !47
|
| 165 |
+
%140 = getelementptr float, ptr addrspace(3) @global_smem, i32 %14, !dbg !47
|
| 166 |
+
%141 = tail call i32 asm sideeffect "@$2 ld.shared.b32 $0, [ $1 + 0 ];", "=r,r,b"(ptr addrspace(3) %140, i1 %139) #4, !dbg !47
|
| 167 |
+
%142 = bitcast i32 %141 to float, !dbg !47
|
| 168 |
+
%143 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %141, i32 1, i32 31), !dbg !47
|
| 169 |
+
%144 = bitcast i32 %143 to float, !dbg !47
|
| 170 |
+
%145 = fadd float %142, %144, !dbg !43
|
| 171 |
+
%146 = and i32 %14, 897, !dbg !47
|
| 172 |
+
%147 = icmp eq i32 %146, 0, !dbg !47
|
| 173 |
+
%148 = bitcast float %145 to <1 x i32>, !dbg !47
|
| 174 |
+
tail call void asm sideeffect "@$2 st.shared.b32 [ $0 + 0 ], $1;", "r,r,b"(ptr addrspace(3) %140, <1 x i32> %148, i1 %147) #4, !dbg !47
|
| 175 |
+
tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !47
|
| 176 |
+
%149 = load i32, ptr addrspace(3) %137, align 8, !dbg !47
|
| 177 |
+
%150 = getelementptr float, ptr addrspace(1) %3, i64 %19, !dbg !48
|
| 178 |
+
%151 = and i32 %14, 64, !dbg !49
|
| 179 |
+
%152 = icmp eq i32 %151, 0, !dbg !49
|
| 180 |
+
%153 = and i1 %152, %17, !dbg !49
|
| 181 |
+
tail call void asm sideeffect "@$2 st.global.b32 [ $1 + 0 ], { $0 };", "r,l,b"(i32 %149, ptr addrspace(1) %150, i1 %153) #4, !dbg !49
|
| 182 |
+
ret void, !dbg !50
|
| 183 |
+
}
|
| 184 |
+
|
| 185 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 186 |
+
declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1
|
| 187 |
+
|
| 188 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 189 |
+
declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
|
| 190 |
+
|
| 191 |
+
; Function Attrs: convergent nocallback nounwind
|
| 192 |
+
declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #2
|
| 193 |
+
|
| 194 |
+
; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
|
| 195 |
+
declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) #3
|
| 196 |
+
|
| 197 |
+
attributes #0 = { nounwind "nvvm.reqntid"="128" }
|
| 198 |
+
attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
|
| 199 |
+
attributes #2 = { convergent nocallback nounwind }
|
| 200 |
+
attributes #3 = { convergent nocallback nounwind memory(inaccessiblemem: readwrite) }
|
| 201 |
+
attributes #4 = { nounwind }
|
| 202 |
+
|
| 203 |
+
!llvm.dbg.cu = !{!0}
|
| 204 |
+
!llvm.module.flags = !{!2, !3}
|
| 205 |
+
|
| 206 |
+
!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
|
| 207 |
+
!1 = !DIFile(filename: "c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a")
|
| 208 |
+
!2 = !{i32 2, !"Debug Info Version", i32 3}
|
| 209 |
+
!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
|
| 210 |
+
!4 = distinct !DISubprogram(name: "triton_red_fused__to_copy_mul_sum_0", linkageName: "triton_red_fused__to_copy_mul_sum_0", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
|
| 211 |
+
!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
|
| 212 |
+
!6 = !{}
|
| 213 |
+
!7 = !DILocation(line: 21, column: 28, scope: !4)
|
| 214 |
+
!8 = !DILocation(line: 21, column: 33, scope: !4)
|
| 215 |
+
!9 = !DILocation(line: 22, column: 44, scope: !4)
|
| 216 |
+
!10 = !DILocation(line: 22, column: 23, scope: !4)
|
| 217 |
+
!11 = !DILocation(line: 23, column: 21, scope: !4)
|
| 218 |
+
!12 = !DILocation(line: 24, column: 37, scope: !4)
|
| 219 |
+
!13 = !DILocation(line: 26, column: 19, scope: !4)
|
| 220 |
+
!14 = !DILocation(line: 27, column: 19, scope: !4)
|
| 221 |
+
!15 = !DILocation(line: 36, column: 36, scope: !4)
|
| 222 |
+
!16 = !DILocation(line: 36, column: 32, scope: !4)
|
| 223 |
+
!17 = !DILocation(line: 36, column: 44, scope: !4)
|
| 224 |
+
!18 = !DILocation(line: 36, column: 26, scope: !4)
|
| 225 |
+
!19 = !DILocation(line: 30, column: 40, scope: !4)
|
| 226 |
+
!20 = !DILocation(line: 39, column: 96, scope: !4)
|
| 227 |
+
!21 = !DILocation(line: 31, column: 31, scope: !4)
|
| 228 |
+
!22 = !DILocation(line: 32, column: 29, scope: !4)
|
| 229 |
+
!23 = !DILocation(line: 36, column: 22, scope: !4)
|
| 230 |
+
!24 = !DILocation(line: 38, column: 22, scope: !4)
|
| 231 |
+
!25 = !DILocation(line: 39, column: 83, scope: !4)
|
| 232 |
+
!26 = !DILocation(line: 39, column: 45, scope: !4)
|
| 233 |
+
!27 = !DILocation(line: 39, column: 39, scope: !4)
|
| 234 |
+
!28 = !DILocation(line: 39, column: 34, scope: !4)
|
| 235 |
+
!29 = !DILocation(line: 39, column: 106, scope: !4)
|
| 236 |
+
!30 = !DILocation(line: 39, column: 113, scope: !4)
|
| 237 |
+
!31 = !DILocation(line: 39, column: 164, scope: !4)
|
| 238 |
+
!32 = !DILocation(line: 40, column: 34, scope: !4)
|
| 239 |
+
!33 = !DILocation(line: 40, column: 96, scope: !4)
|
| 240 |
+
!34 = !DILocation(line: 40, column: 164, scope: !4)
|
| 241 |
+
!35 = !DILocation(line: 42, column: 35, scope: !4)
|
| 242 |
+
!36 = !DILocation(line: 42, column: 85, scope: !4)
|
| 243 |
+
!37 = !DILocation(line: 43, column: 22, scope: !4)
|
| 244 |
+
!38 = !DILocation(line: 45, column: 22, scope: !4)
|
| 245 |
+
!39 = !DILocation(line: 47, column: 37, scope: !4)
|
| 246 |
+
!40 = !DILocation(line: 49, column: 25, scope: !4)
|
| 247 |
+
!41 = !DILocation(line: 50, column: 36, scope: !4)
|
| 248 |
+
!42 = !DILocation(line: 50, column: 50, scope: !4)
|
| 249 |
+
!43 = !DILocation(line: 261, column: 15, scope: !44, inlinedAt: !46)
|
| 250 |
+
!44 = distinct !DILexicalBlockFile(scope: !4, file: !45, discriminator: 0)
|
| 251 |
+
!45 = !DIFile(filename: "standard.py", directory: "/workspace/specforge/lib/python3.11/site-packages/triton/language")
|
| 252 |
+
!46 = !DILocation(line: 51, column: 27, scope: !4)
|
| 253 |
+
!47 = !DILocation(line: 291, column: 36, scope: !44, inlinedAt: !46)
|
| 254 |
+
!48 = !DILocation(line: 52, column: 25, scope: !4)
|
| 255 |
+
!49 = !DILocation(line: 52, column: 37, scope: !4)
|
| 256 |
+
!50 = !DILocation(line: 52, column: 4, scope: !4)
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ptx
ADDED
|
@@ -0,0 +1,688 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
//
|
| 2 |
+
// Generated by LLVM NVPTX Back-End
|
| 3 |
+
//
|
| 4 |
+
|
| 5 |
+
.version 8.7
|
| 6 |
+
.target sm_90a
|
| 7 |
+
.address_size 64
|
| 8 |
+
|
| 9 |
+
// .globl triton_red_fused__to_copy_mul_sum_0 // -- Begin function triton_red_fused__to_copy_mul_sum_0
|
| 10 |
+
.extern .shared .align 16 .b8 global_smem[];
|
| 11 |
+
// @triton_red_fused__to_copy_mul_sum_0
|
| 12 |
+
.visible .entry triton_red_fused__to_copy_mul_sum_0(
|
| 13 |
+
.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_mul_sum_0_param_0,
|
| 14 |
+
.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_mul_sum_0_param_1,
|
| 15 |
+
.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_mul_sum_0_param_2,
|
| 16 |
+
.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_mul_sum_0_param_3,
|
| 17 |
+
.param .u64 triton_red_fused__to_copy_mul_sum_0_param_4,
|
| 18 |
+
.param .u64 triton_red_fused__to_copy_mul_sum_0_param_5,
|
| 19 |
+
.param .u64 triton_red_fused__to_copy_mul_sum_0_param_6,
|
| 20 |
+
.param .u32 triton_red_fused__to_copy_mul_sum_0_param_7,
|
| 21 |
+
.param .u32 triton_red_fused__to_copy_mul_sum_0_param_8,
|
| 22 |
+
.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_mul_sum_0_param_9,
|
| 23 |
+
.param .u64 .ptr .global .align 1 triton_red_fused__to_copy_mul_sum_0_param_10
|
| 24 |
+
)
|
| 25 |
+
.reqntid 128
|
| 26 |
+
{
|
| 27 |
+
.reg .pred %p<43>;
|
| 28 |
+
.reg .b16 %rs<25>;
|
| 29 |
+
.reg .b32 %r<119>;
|
| 30 |
+
.reg .b64 %rd<137>;
|
| 31 |
+
.loc 1 18 0 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:18:0
|
| 32 |
+
$L__func_begin0:
|
| 33 |
+
.loc 1 18 0 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:18:0
|
| 34 |
+
|
| 35 |
+
// %bb.0:
|
| 36 |
+
ld.param.b32 %r17, [triton_red_fused__to_copy_mul_sum_0_param_8];
|
| 37 |
+
ld.param.b64 %rd46, [triton_red_fused__to_copy_mul_sum_0_param_4];
|
| 38 |
+
$L__tmp0:
|
| 39 |
+
.loc 1 21 28 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:21:28
|
| 40 |
+
mov.u32 %r18, %ctaid.x;
|
| 41 |
+
.loc 1 21 33 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:21:33
|
| 42 |
+
shl.b32 %r19, %r18, 6;
|
| 43 |
+
.loc 1 22 44 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:22:44
|
| 44 |
+
mov.u32 %r1, %tid.x;
|
| 45 |
+
and.b32 %r2, %r1, 63;
|
| 46 |
+
.loc 1 22 23 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:22:23
|
| 47 |
+
or.b32 %r20, %r19, %r2;
|
| 48 |
+
.loc 1 26 19 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:26:19
|
| 49 |
+
cvt.s64.s32 %rd1, %r20;
|
| 50 |
+
or.b64 %rd50, %rd1, %rd46;
|
| 51 |
+
and.b64 %rd51, %rd50, -4294967296;
|
| 52 |
+
setp.ne.b64 %p5, %rd51, 0;
|
| 53 |
+
cvt.u32.u64 %r116, %rd1;
|
| 54 |
+
@%p5 bra $L__BB0_2;
|
| 55 |
+
bra.uni $L__BB0_1;
|
| 56 |
+
$L__BB0_2:
|
| 57 |
+
div.s64 %rd130, %rd1, %rd46;
|
| 58 |
+
bra.uni $L__BB0_3;
|
| 59 |
+
$L__BB0_1:
|
| 60 |
+
cvt.u32.u64 %r21, %rd46;
|
| 61 |
+
div.u32 %r23, %r116, %r21;
|
| 62 |
+
cvt.u64.u32 %rd130, %r23;
|
| 63 |
+
$L__BB0_3:
|
| 64 |
+
.loc 1 0 19 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:0:19
|
| 65 |
+
ld.param.b32 %r16, [triton_red_fused__to_copy_mul_sum_0_param_7];
|
| 66 |
+
ld.param.b64 %rd45, [triton_red_fused__to_copy_mul_sum_0_param_3];
|
| 67 |
+
bfe.u32 %r3, %r1, 6, 1;
|
| 68 |
+
.loc 1 30 40 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:30:40
|
| 69 |
+
setp.lt.s32 %p6, %r17, 1;
|
| 70 |
+
mov.b32 %r118, 0;
|
| 71 |
+
@%p6 bra $L__BB0_19;
|
| 72 |
+
// %bb.4: // %.lr.ph.preheader
|
| 73 |
+
.loc 1 0 40 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:0:40
|
| 74 |
+
ld.param.b64 %rd48, [triton_red_fused__to_copy_mul_sum_0_param_6];
|
| 75 |
+
ld.param.b64 %rd47, [triton_red_fused__to_copy_mul_sum_0_param_5];
|
| 76 |
+
ld.param.b64 %rd44, [triton_red_fused__to_copy_mul_sum_0_param_2];
|
| 77 |
+
ld.param.b64 %rd43, [triton_red_fused__to_copy_mul_sum_0_param_1];
|
| 78 |
+
ld.param.b64 %rd42, [triton_red_fused__to_copy_mul_sum_0_param_0];
|
| 79 |
+
mul.lo.s64 %rd52, %rd130, %rd46;
|
| 80 |
+
sub.s64 %rd6, %rd1, %rd52;
|
| 81 |
+
mul.lo.s64 %rd7, %rd48, %rd47;
|
| 82 |
+
add.s64 %rd53, %rd7, 31;
|
| 83 |
+
shr.s64 %rd54, %rd53, 63;
|
| 84 |
+
shr.u64 %rd55, %rd54, 59;
|
| 85 |
+
add.s64 %rd56, %rd53, %rd55;
|
| 86 |
+
shr.s64 %rd57, %rd56, 5;
|
| 87 |
+
mul.lo.s64 %rd8, %rd130, %rd57;
|
| 88 |
+
.loc 1 23 21 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:23:21
|
| 89 |
+
setp.lt.s32 %p1, %r116, %r16;
|
| 90 |
+
mov.b32 %r27, 0f00000000;
|
| 91 |
+
mov.b64 %rd131, {%r27, %r27};
|
| 92 |
+
mov.b32 %r117, 0;
|
| 93 |
+
mov.b64 %rd132, %rd131;
|
| 94 |
+
bra.uni $L__BB0_5;
|
| 95 |
+
$L__BB0_16: // in Loop: Header=BB0_5 Depth=1
|
| 96 |
+
.loc 1 39 83 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:83
|
| 97 |
+
rem.s64 %rd136, %rd23, %rd7;
|
| 98 |
+
$L__BB0_17: // in Loop: Header=BB0_5 Depth=1
|
| 99 |
+
.loc 1 38 22 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:38:22
|
| 100 |
+
setp.lt.s64 %p23, %rd20, %rd7;
|
| 101 |
+
setp.lt.s64 %p24, %rd21, %rd7;
|
| 102 |
+
setp.lt.s64 %p25, %rd22, %rd7;
|
| 103 |
+
setp.lt.s64 %p26, %rd23, %rd7;
|
| 104 |
+
.loc 1 32 29 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:32:29
|
| 105 |
+
setp.lt.s32 %p27, %r9, %r17;
|
| 106 |
+
setp.lt.s32 %p28, %r10, %r17;
|
| 107 |
+
setp.lt.s32 %p29, %r11, %r17;
|
| 108 |
+
setp.lt.s32 %p30, %r12, %r17;
|
| 109 |
+
.loc 1 39 39 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:39
|
| 110 |
+
mad.lo.s64 %rd108, %rd133, %rd46, %rd6;
|
| 111 |
+
mad.lo.s64 %rd109, %rd134, %rd46, %rd6;
|
| 112 |
+
mad.lo.s64 %rd110, %rd135, %rd46, %rd6;
|
| 113 |
+
mad.lo.s64 %rd111, %rd136, %rd46, %rd6;
|
| 114 |
+
.loc 1 39 34 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:34
|
| 115 |
+
shl.b64 %rd112, %rd108, 1;
|
| 116 |
+
add.s64 %rd73, %rd42, %rd112;
|
| 117 |
+
shl.b64 %rd113, %rd109, 1;
|
| 118 |
+
add.s64 %rd76, %rd42, %rd113;
|
| 119 |
+
shl.b64 %rd114, %rd110, 1;
|
| 120 |
+
add.s64 %rd79, %rd42, %rd114;
|
| 121 |
+
shl.b64 %rd115, %rd111, 1;
|
| 122 |
+
add.s64 %rd82, %rd42, %rd115;
|
| 123 |
+
.loc 1 39 106 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:106
|
| 124 |
+
and.pred %p32, %p30, %p26;
|
| 125 |
+
and.pred %p33, %p29, %p25;
|
| 126 |
+
and.pred %p34, %p28, %p24;
|
| 127 |
+
and.pred %p35, %p27, %p23;
|
| 128 |
+
.loc 1 39 113 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:113
|
| 129 |
+
and.pred %p11, %p1, %p35;
|
| 130 |
+
and.pred %p12, %p1, %p34;
|
| 131 |
+
and.pred %p13, %p1, %p33;
|
| 132 |
+
and.pred %p14, %p1, %p32;
|
| 133 |
+
mov.b16 %rs2, 0;
|
| 134 |
+
.loc 1 39 96 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:96
|
| 135 |
+
// begin inline asm
|
| 136 |
+
mov.u16 %rs1, %rs2;
|
| 137 |
+
@%p11 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs1 }, [ %rd73 + 0 ], %rd60;
|
| 138 |
+
// end inline asm
|
| 139 |
+
// begin inline asm
|
| 140 |
+
mov.u64 %rd75, 0x0;
|
| 141 |
+
createpolicy.fractional.L2::evict_last.b64 %rd75, 1.0;
|
| 142 |
+
// end inline asm
|
| 143 |
+
// begin inline asm
|
| 144 |
+
mov.u16 %rs3, %rs2;
|
| 145 |
+
@%p12 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs3 }, [ %rd76 + 0 ], %rd75;
|
| 146 |
+
// end inline asm
|
| 147 |
+
// begin inline asm
|
| 148 |
+
mov.u64 %rd78, 0x0;
|
| 149 |
+
createpolicy.fractional.L2::evict_last.b64 %rd78, 1.0;
|
| 150 |
+
// end inline asm
|
| 151 |
+
// begin inline asm
|
| 152 |
+
mov.u16 %rs5, %rs2;
|
| 153 |
+
@%p13 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs5 }, [ %rd79 + 0 ], %rd78;
|
| 154 |
+
// end inline asm
|
| 155 |
+
// begin inline asm
|
| 156 |
+
mov.u64 %rd81, 0x0;
|
| 157 |
+
createpolicy.fractional.L2::evict_last.b64 %rd81, 1.0;
|
| 158 |
+
// end inline asm
|
| 159 |
+
// begin inline asm
|
| 160 |
+
mov.u16 %rs7, %rs2;
|
| 161 |
+
@%p14 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs7 }, [ %rd82 + 0 ], %rd81;
|
| 162 |
+
// end inline asm
|
| 163 |
+
mov.b32 %r49, {%rs5, %rs7};
|
| 164 |
+
mov.b32 %r50, {%rs1, %rs3};
|
| 165 |
+
.loc 1 39 164 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:164
|
| 166 |
+
mov.b32 {%rs17, %rs18}, %r50;
|
| 167 |
+
cvt.f32.bf16 %r51, %rs18;
|
| 168 |
+
cvt.f32.bf16 %r52, %rs17;
|
| 169 |
+
mov.b32 {%rs19, %rs20}, %r49;
|
| 170 |
+
cvt.f32.bf16 %r53, %rs20;
|
| 171 |
+
cvt.f32.bf16 %r54, %rs19;
|
| 172 |
+
.loc 1 40 34 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:40:34
|
| 173 |
+
add.s64 %rd85, %rd43, %rd112;
|
| 174 |
+
add.s64 %rd88, %rd43, %rd113;
|
| 175 |
+
add.s64 %rd91, %rd43, %rd114;
|
| 176 |
+
add.s64 %rd94, %rd43, %rd115;
|
| 177 |
+
.loc 1 40 96 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:40:96
|
| 178 |
+
// begin inline asm
|
| 179 |
+
mov.u64 %rd84, 0x0;
|
| 180 |
+
createpolicy.fractional.L2::evict_last.b64 %rd84, 1.0;
|
| 181 |
+
// end inline asm
|
| 182 |
+
// begin inline asm
|
| 183 |
+
mov.u16 %rs9, %rs2;
|
| 184 |
+
@%p11 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs9 }, [ %rd85 + 0 ], %rd84;
|
| 185 |
+
// end inline asm
|
| 186 |
+
// begin inline asm
|
| 187 |
+
mov.u64 %rd87, 0x0;
|
| 188 |
+
createpolicy.fractional.L2::evict_last.b64 %rd87, 1.0;
|
| 189 |
+
// end inline asm
|
| 190 |
+
// begin inline asm
|
| 191 |
+
mov.u16 %rs11, %rs2;
|
| 192 |
+
@%p12 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs11 }, [ %rd88 + 0 ], %rd87;
|
| 193 |
+
// end inline asm
|
| 194 |
+
// begin inline asm
|
| 195 |
+
mov.u64 %rd90, 0x0;
|
| 196 |
+
createpolicy.fractional.L2::evict_last.b64 %rd90, 1.0;
|
| 197 |
+
// end inline asm
|
| 198 |
+
// begin inline asm
|
| 199 |
+
mov.u16 %rs13, %rs2;
|
| 200 |
+
@%p13 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs13 }, [ %rd91 + 0 ], %rd90;
|
| 201 |
+
// end inline asm
|
| 202 |
+
// begin inline asm
|
| 203 |
+
mov.u64 %rd93, 0x0;
|
| 204 |
+
createpolicy.fractional.L2::evict_last.b64 %rd93, 1.0;
|
| 205 |
+
// end inline asm
|
| 206 |
+
// begin inline asm
|
| 207 |
+
mov.u16 %rs15, %rs2;
|
| 208 |
+
@%p14 ld.global.L1::evict_last.L2::cache_hint.b16 { %rs15 }, [ %rd94 + 0 ], %rd93;
|
| 209 |
+
// end inline asm
|
| 210 |
+
mov.b32 %r55, {%rs9, %rs11};
|
| 211 |
+
mov.b32 %r56, {%rs13, %rs15};
|
| 212 |
+
.loc 1 40 164 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:40:164
|
| 213 |
+
mov.b32 {%rs21, %rs22}, %r56;
|
| 214 |
+
cvt.f32.bf16 %r57, %rs21;
|
| 215 |
+
cvt.f32.bf16 %r58, %rs22;
|
| 216 |
+
mov.b32 {%rs23, %rs24}, %r55;
|
| 217 |
+
cvt.f32.bf16 %r59, %rs23;
|
| 218 |
+
cvt.f32.bf16 %r60, %rs24;
|
| 219 |
+
.loc 1 42 35 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:42:35
|
| 220 |
+
shl.b64 %rd116, %rd133, 2;
|
| 221 |
+
add.s64 %rd97, %rd44, %rd116;
|
| 222 |
+
shl.b64 %rd117, %rd134, 2;
|
| 223 |
+
add.s64 %rd100, %rd44, %rd117;
|
| 224 |
+
shl.b64 %rd118, %rd135, 2;
|
| 225 |
+
add.s64 %rd103, %rd44, %rd118;
|
| 226 |
+
shl.b64 %rd119, %rd136, 2;
|
| 227 |
+
add.s64 %rd106, %rd44, %rd119;
|
| 228 |
+
.loc 1 42 85 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:42:85
|
| 229 |
+
// begin inline asm
|
| 230 |
+
mov.u64 %rd96, 0x0;
|
| 231 |
+
createpolicy.fractional.L2::evict_last.b64 %rd96, 1.0;
|
| 232 |
+
// end inline asm
|
| 233 |
+
mov.b32 %r41, 0;
|
| 234 |
+
// begin inline asm
|
| 235 |
+
mov.u32 %r40, %r41;
|
| 236 |
+
@%p11 ld.global.L1::evict_last.L2::cache_hint.b32 { %r40 }, [ %rd97 + 0 ], %rd96;
|
| 237 |
+
// end inline asm
|
| 238 |
+
// begin inline asm
|
| 239 |
+
mov.u64 %rd99, 0x0;
|
| 240 |
+
createpolicy.fractional.L2::evict_last.b64 %rd99, 1.0;
|
| 241 |
+
// end inline asm
|
| 242 |
+
// begin inline asm
|
| 243 |
+
mov.u32 %r42, %r41;
|
| 244 |
+
@%p12 ld.global.L1::evict_last.L2::cache_hint.b32 { %r42 }, [ %rd100 + 0 ], %rd99;
|
| 245 |
+
// end inline asm
|
| 246 |
+
// begin inline asm
|
| 247 |
+
mov.u64 %rd102, 0x0;
|
| 248 |
+
createpolicy.fractional.L2::evict_last.b64 %rd102, 1.0;
|
| 249 |
+
// end inline asm
|
| 250 |
+
// begin inline asm
|
| 251 |
+
mov.u32 %r44, %r41;
|
| 252 |
+
@%p13 ld.global.L1::evict_last.L2::cache_hint.b32 { %r44 }, [ %rd103 + 0 ], %rd102;
|
| 253 |
+
// end inline asm
|
| 254 |
+
// begin inline asm
|
| 255 |
+
mov.u64 %rd105, 0x0;
|
| 256 |
+
createpolicy.fractional.L2::evict_last.b64 %rd105, 1.0;
|
| 257 |
+
// end inline asm
|
| 258 |
+
// begin inline asm
|
| 259 |
+
mov.u32 %r46, %r41;
|
| 260 |
+
@%p14 ld.global.L1::evict_last.L2::cache_hint.b32 { %r46 }, [ %rd106 + 0 ], %rd105;
|
| 261 |
+
// end inline asm
|
| 262 |
+
cvt.u64.u32 %rd120, %r44;
|
| 263 |
+
cvt.u64.u32 %rd121, %r46;
|
| 264 |
+
shl.b64 %rd122, %rd121, 32;
|
| 265 |
+
or.b64 %rd123, %rd120, %rd122;
|
| 266 |
+
cvt.u64.u32 %rd124, %r40;
|
| 267 |
+
cvt.u64.u32 %rd125, %r42;
|
| 268 |
+
shl.b64 %rd126, %rd125, 32;
|
| 269 |
+
or.b64 %rd127, %rd124, %rd126;
|
| 270 |
+
.loc 1 43 22 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:43:22
|
| 271 |
+
mov.b64 {%r61, %r62}, %rd127;
|
| 272 |
+
mul.f32 %r63, %r60, %r62;
|
| 273 |
+
mul.f32 %r64, %r59, %r61;
|
| 274 |
+
mov.b64 {%r65, %r66}, %rd123;
|
| 275 |
+
mul.f32 %r67, %r58, %r66;
|
| 276 |
+
mul.f32 %r68, %r57, %r65;
|
| 277 |
+
.loc 1 45 22 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:45:22
|
| 278 |
+
mul.f32 %r69, %r68, %r54;
|
| 279 |
+
mul.f32 %r70, %r67, %r53;
|
| 280 |
+
mul.f32 %r71, %r64, %r52;
|
| 281 |
+
mul.f32 %r72, %r63, %r51;
|
| 282 |
+
.loc 1 47 37 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:47:37
|
| 283 |
+
selp.f32 %r73, %r72, 0f00000000, %p24;
|
| 284 |
+
selp.f32 %r74, %r71, 0f00000000, %p23;
|
| 285 |
+
selp.f32 %r75, %r70, 0f00000000, %p26;
|
| 286 |
+
selp.f32 %r76, %r69, 0f00000000, %p25;
|
| 287 |
+
.loc 1 49 25 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:49:25
|
| 288 |
+
mov.b64 {%r77, %r78}, %rd132;
|
| 289 |
+
add.f32 %r79, %r77, %r76;
|
| 290 |
+
add.f32 %r80, %r78, %r75;
|
| 291 |
+
mov.b64 {%r81, %r82}, %rd131;
|
| 292 |
+
add.f32 %r83, %r81, %r74;
|
| 293 |
+
add.f32 %r84, %r82, %r73;
|
| 294 |
+
.loc 1 50 50 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:50:50
|
| 295 |
+
selp.f32 %r85, %r84, %r82, %p28;
|
| 296 |
+
selp.f32 %r86, %r85, %r82, %p1;
|
| 297 |
+
selp.f32 %r87, %r83, %r81, %p27;
|
| 298 |
+
selp.f32 %r88, %r87, %r81, %p1;
|
| 299 |
+
mov.b64 %rd131, {%r88, %r86};
|
| 300 |
+
selp.f32 %r89, %r80, %r78, %p30;
|
| 301 |
+
selp.f32 %r90, %r89, %r78, %p1;
|
| 302 |
+
selp.f32 %r91, %r79, %r77, %p29;
|
| 303 |
+
selp.f32 %r92, %r91, %r77, %p1;
|
| 304 |
+
mov.b64 %rd132, {%r92, %r90};
|
| 305 |
+
.loc 1 30 40 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:30:40
|
| 306 |
+
add.s32 %r117, %r117, 8;
|
| 307 |
+
setp.lt.s32 %p36, %r117, %r17;
|
| 308 |
+
@%p36 bra $L__BB0_5;
|
| 309 |
+
bra.uni $L__BB0_18;
|
| 310 |
+
$L__BB0_5: // %.lr.ph
|
| 311 |
+
// =>This Inner Loop Header: Depth=1
|
| 312 |
+
.loc 1 39 96 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:96
|
| 313 |
+
// begin inline asm
|
| 314 |
+
mov.u64 %rd60, 0x0;
|
| 315 |
+
createpolicy.fractional.L2::evict_last.b64 %rd60, 1.0;
|
| 316 |
+
// end inline asm
|
| 317 |
+
.loc 1 31 31 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:31:31
|
| 318 |
+
add.s32 %r9, %r3, %r117;
|
| 319 |
+
add.s32 %r10, %r9, 2;
|
| 320 |
+
.loc 1 36 22 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:36:22
|
| 321 |
+
cvt.s64.s32 %rd61, %r9;
|
| 322 |
+
cvt.s64.s32 %rd62, %r10;
|
| 323 |
+
add.s64 %rd21, %rd8, %rd62;
|
| 324 |
+
add.s64 %rd20, %rd8, %rd61;
|
| 325 |
+
.loc 1 39 83 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:83
|
| 326 |
+
or.b64 %rd65, %rd20, %rd7;
|
| 327 |
+
and.b64 %rd66, %rd65, -4294967296;
|
| 328 |
+
setp.ne.b64 %p7, %rd66, 0;
|
| 329 |
+
@%p7 bra $L__BB0_7;
|
| 330 |
+
bra.uni $L__BB0_6;
|
| 331 |
+
$L__BB0_7: // in Loop: Header=BB0_5 Depth=1
|
| 332 |
+
rem.s64 %rd133, %rd20, %rd7;
|
| 333 |
+
bra.uni $L__BB0_8;
|
| 334 |
+
$L__BB0_6: // in Loop: Header=BB0_5 Depth=1
|
| 335 |
+
cvt.u32.u64 %r28, %rd7;
|
| 336 |
+
cvt.u32.u64 %r29, %rd20;
|
| 337 |
+
rem.u32 %r30, %r29, %r28;
|
| 338 |
+
cvt.u64.u32 %rd133, %r30;
|
| 339 |
+
$L__BB0_8: // in Loop: Header=BB0_5 Depth=1
|
| 340 |
+
.loc 1 0 0 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:0
|
| 341 |
+
add.s32 %r11, %r9, 4;
|
| 342 |
+
cvt.s64.s32 %rd63, %r11;
|
| 343 |
+
add.s64 %rd22, %rd8, %rd63;
|
| 344 |
+
.loc 1 39 83 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:83
|
| 345 |
+
or.b64 %rd67, %rd21, %rd7;
|
| 346 |
+
and.b64 %rd68, %rd67, -4294967296;
|
| 347 |
+
setp.ne.b64 %p8, %rd68, 0;
|
| 348 |
+
@%p8 bra $L__BB0_10;
|
| 349 |
+
bra.uni $L__BB0_9;
|
| 350 |
+
$L__BB0_10: // in Loop: Header=BB0_5 Depth=1
|
| 351 |
+
rem.s64 %rd134, %rd21, %rd7;
|
| 352 |
+
bra.uni $L__BB0_11;
|
| 353 |
+
$L__BB0_9: // in Loop: Header=BB0_5 Depth=1
|
| 354 |
+
cvt.u32.u64 %r31, %rd7;
|
| 355 |
+
cvt.u32.u64 %r32, %rd21;
|
| 356 |
+
rem.u32 %r33, %r32, %r31;
|
| 357 |
+
cvt.u64.u32 %rd134, %r33;
|
| 358 |
+
$L__BB0_11: // in Loop: Header=BB0_5 Depth=1
|
| 359 |
+
.loc 1 0 0 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:0
|
| 360 |
+
add.s32 %r12, %r9, 6;
|
| 361 |
+
cvt.s64.s32 %rd64, %r12;
|
| 362 |
+
add.s64 %rd23, %rd8, %rd64;
|
| 363 |
+
.loc 1 39 83 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:39:83
|
| 364 |
+
or.b64 %rd69, %rd22, %rd7;
|
| 365 |
+
and.b64 %rd70, %rd69, -4294967296;
|
| 366 |
+
setp.ne.b64 %p9, %rd70, 0;
|
| 367 |
+
@%p9 bra $L__BB0_13;
|
| 368 |
+
bra.uni $L__BB0_12;
|
| 369 |
+
$L__BB0_13: // in Loop: Header=BB0_5 Depth=1
|
| 370 |
+
rem.s64 %rd135, %rd22, %rd7;
|
| 371 |
+
bra.uni $L__BB0_14;
|
| 372 |
+
$L__BB0_12: // in Loop: Header=BB0_5 Depth=1
|
| 373 |
+
cvt.u32.u64 %r34, %rd7;
|
| 374 |
+
cvt.u32.u64 %r35, %rd22;
|
| 375 |
+
rem.u32 %r36, %r35, %r34;
|
| 376 |
+
cvt.u64.u32 %rd135, %r36;
|
| 377 |
+
$L__BB0_14: // in Loop: Header=BB0_5 Depth=1
|
| 378 |
+
or.b64 %rd71, %rd23, %rd7;
|
| 379 |
+
and.b64 %rd72, %rd71, -4294967296;
|
| 380 |
+
setp.ne.b64 %p10, %rd72, 0;
|
| 381 |
+
@%p10 bra $L__BB0_16;
|
| 382 |
+
// %bb.15: // in Loop: Header=BB0_5 Depth=1
|
| 383 |
+
cvt.u32.u64 %r37, %rd7;
|
| 384 |
+
cvt.u32.u64 %r38, %rd23;
|
| 385 |
+
rem.u32 %r39, %r38, %r37;
|
| 386 |
+
cvt.u64.u32 %rd136, %r39;
|
| 387 |
+
bra.uni $L__BB0_17;
|
| 388 |
+
$L__BB0_18: // %._crit_edge.loopexit
|
| 389 |
+
$L__tmp1:
|
| 390 |
+
.loc 2 261 15 // standard.py:261:15 @[ c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:51:27 ]
|
| 391 |
+
mov.b64 {%r93, %r94}, %rd131;
|
| 392 |
+
add.f32 %r95, %r93, %r94;
|
| 393 |
+
mov.b64 {%r96, %r97}, %rd132;
|
| 394 |
+
add.f32 %r98, %r96, %r95;
|
| 395 |
+
add.f32 %r118, %r97, %r98;
|
| 396 |
+
$L__tmp2:
|
| 397 |
+
$L__BB0_19: // %._crit_edge
|
| 398 |
+
.loc 1 23 21 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:23:21
|
| 399 |
+
setp.lt.s32 %p41, %r116, %r16;
|
| 400 |
+
$L__tmp3:
|
| 401 |
+
.loc 2 291 36 // standard.py:291:36 @[ c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:51:27 ]
|
| 402 |
+
shl.b32 %r107, %r2, 3;
|
| 403 |
+
mov.b32 %r108, global_smem;
|
| 404 |
+
add.s32 %r109, %r108, %r107;
|
| 405 |
+
shl.b32 %r110, %r3, 2;
|
| 406 |
+
add.s32 %r99, %r109, %r110;
|
| 407 |
+
mov.pred %p37, -1;
|
| 408 |
+
// begin inline asm
|
| 409 |
+
@%p37 st.shared.b32 [ %r99 + 0 ], %r118;
|
| 410 |
+
// end inline asm
|
| 411 |
+
bar.sync 0;
|
| 412 |
+
setp.lt.u32 %p38, %r1, 128;
|
| 413 |
+
shl.b32 %r111, %r1, 2;
|
| 414 |
+
add.s32 %r102, %r108, %r111;
|
| 415 |
+
// begin inline asm
|
| 416 |
+
@%p38 ld.shared.b32 %r101, [ %r102 + 0 ];
|
| 417 |
+
// end inline asm
|
| 418 |
+
shfl.sync.bfly.b32 %r112, %r101, 1, 31, -1;
|
| 419 |
+
.loc 2 261 15 // standard.py:261:15 @[ c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:51:27 ]
|
| 420 |
+
add.f32 %r104, %r101, %r112;
|
| 421 |
+
.loc 2 291 36 // standard.py:291:36 @[ c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:51:27 ]
|
| 422 |
+
and.b32 %r113, %r1, 897;
|
| 423 |
+
setp.eq.b32 %p39, %r113, 0;
|
| 424 |
+
// begin inline asm
|
| 425 |
+
@%p39 st.shared.b32 [ %r102 + 0 ], %r104;
|
| 426 |
+
// end inline asm
|
| 427 |
+
bar.sync 0;
|
| 428 |
+
ld.shared.b32 %r105, [%r109];
|
| 429 |
+
$L__tmp4:
|
| 430 |
+
.loc 1 52 25 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:52:25
|
| 431 |
+
shl.b64 %rd129, %rd1, 2;
|
| 432 |
+
add.s64 %rd128, %rd45, %rd129;
|
| 433 |
+
.loc 1 52 37 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:52:37
|
| 434 |
+
and.b32 %r114, %r1, 64;
|
| 435 |
+
setp.eq.b32 %p42, %r114, 0;
|
| 436 |
+
and.pred %p40, %p42, %p41;
|
| 437 |
+
// begin inline asm
|
| 438 |
+
@%p40 st.global.b32 [ %rd128 + 0 ], { %r105 };
|
| 439 |
+
// end inline asm
|
| 440 |
+
.loc 1 52 4 // c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py:52:4
|
| 441 |
+
ret;
|
| 442 |
+
$L__tmp5:
|
| 443 |
+
$L__func_end0:
|
| 444 |
+
// -- End function
|
| 445 |
+
}
|
| 446 |
+
.file 1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py"
|
| 447 |
+
.file 2 "/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py"
|
| 448 |
+
.section .debug_abbrev
|
| 449 |
+
{
|
| 450 |
+
.b8 1 // Abbreviation Code
|
| 451 |
+
.b8 17 // DW_TAG_compile_unit
|
| 452 |
+
.b8 1 // DW_CHILDREN_yes
|
| 453 |
+
.b8 37 // DW_AT_producer
|
| 454 |
+
.b8 8 // DW_FORM_string
|
| 455 |
+
.b8 19 // DW_AT_language
|
| 456 |
+
.b8 5 // DW_FORM_data2
|
| 457 |
+
.b8 3 // DW_AT_name
|
| 458 |
+
.b8 8 // DW_FORM_string
|
| 459 |
+
.b8 16 // DW_AT_stmt_list
|
| 460 |
+
.b8 6 // DW_FORM_data4
|
| 461 |
+
.b8 27 // DW_AT_comp_dir
|
| 462 |
+
.b8 8 // DW_FORM_string
|
| 463 |
+
.b8 0 // EOM(1)
|
| 464 |
+
.b8 0 // EOM(2)
|
| 465 |
+
.b8 2 // Abbreviation Code
|
| 466 |
+
.b8 46 // DW_TAG_subprogram
|
| 467 |
+
.b8 0 // DW_CHILDREN_no
|
| 468 |
+
.b8 3 // DW_AT_name
|
| 469 |
+
.b8 8 // DW_FORM_string
|
| 470 |
+
.b8 32 // DW_AT_inline
|
| 471 |
+
.b8 11 // DW_FORM_data1
|
| 472 |
+
.b8 0 // EOM(1)
|
| 473 |
+
.b8 0 // EOM(2)
|
| 474 |
+
.b8 3 // Abbreviation Code
|
| 475 |
+
.b8 46 // DW_TAG_subprogram
|
| 476 |
+
.b8 1 // DW_CHILDREN_yes
|
| 477 |
+
.b8 17 // DW_AT_low_pc
|
| 478 |
+
.b8 1 // DW_FORM_addr
|
| 479 |
+
.b8 18 // DW_AT_high_pc
|
| 480 |
+
.b8 1 // DW_FORM_addr
|
| 481 |
+
.b8 49 // DW_AT_abstract_origin
|
| 482 |
+
.b8 19 // DW_FORM_ref4
|
| 483 |
+
.b8 0 // EOM(1)
|
| 484 |
+
.b8 0 // EOM(2)
|
| 485 |
+
.b8 4 // Abbreviation Code
|
| 486 |
+
.b8 29 // DW_TAG_inlined_subroutine
|
| 487 |
+
.b8 0 // DW_CHILDREN_no
|
| 488 |
+
.b8 49 // DW_AT_abstract_origin
|
| 489 |
+
.b8 19 // DW_FORM_ref4
|
| 490 |
+
.b8 17 // DW_AT_low_pc
|
| 491 |
+
.b8 1 // DW_FORM_addr
|
| 492 |
+
.b8 18 // DW_AT_high_pc
|
| 493 |
+
.b8 1 // DW_FORM_addr
|
| 494 |
+
.b8 88 // DW_AT_call_file
|
| 495 |
+
.b8 11 // DW_FORM_data1
|
| 496 |
+
.b8 89 // DW_AT_call_line
|
| 497 |
+
.b8 11 // DW_FORM_data1
|
| 498 |
+
.b8 87 // DW_AT_call_column
|
| 499 |
+
.b8 11 // DW_FORM_data1
|
| 500 |
+
.b8 0 // EOM(1)
|
| 501 |
+
.b8 0 // EOM(2)
|
| 502 |
+
.b8 0 // EOM(3)
|
| 503 |
+
}
|
| 504 |
+
.section .debug_info
|
| 505 |
+
{
|
| 506 |
+
.b32 220 // Length of Unit
|
| 507 |
+
.b8 2 // DWARF version number
|
| 508 |
+
.b8 0
|
| 509 |
+
.b32 .debug_abbrev // Offset Into Abbrev. Section
|
| 510 |
+
.b8 8 // Address Size (in bytes)
|
| 511 |
+
.b8 1 // Abbrev [1] 0xb:0xd5 DW_TAG_compile_unit
|
| 512 |
+
.b8 116 // DW_AT_producer
|
| 513 |
+
.b8 114
|
| 514 |
+
.b8 105
|
| 515 |
+
.b8 116
|
| 516 |
+
.b8 111
|
| 517 |
+
.b8 110
|
| 518 |
+
.b8 0
|
| 519 |
+
.b8 2 // DW_AT_language
|
| 520 |
+
.b8 0
|
| 521 |
+
.b8 99 // DW_AT_name
|
| 522 |
+
.b8 53
|
| 523 |
+
.b8 97
|
| 524 |
+
.b8 106
|
| 525 |
+
.b8 52
|
| 526 |
+
.b8 122
|
| 527 |
+
.b8 117
|
| 528 |
+
.b8 114
|
| 529 |
+
.b8 107
|
| 530 |
+
.b8 108
|
| 531 |
+
.b8 99
|
| 532 |
+
.b8 52
|
| 533 |
+
.b8 106
|
| 534 |
+
.b8 104
|
| 535 |
+
.b8 118
|
| 536 |
+
.b8 113
|
| 537 |
+
.b8 112
|
| 538 |
+
.b8 110
|
| 539 |
+
.b8 102
|
| 540 |
+
.b8 103
|
| 541 |
+
.b8 104
|
| 542 |
+
.b8 105
|
| 543 |
+
.b8 51
|
| 544 |
+
.b8 118
|
| 545 |
+
.b8 122
|
| 546 |
+
.b8 54
|
| 547 |
+
.b8 107
|
| 548 |
+
.b8 104
|
| 549 |
+
.b8 106
|
| 550 |
+
.b8 119
|
| 551 |
+
.b8 107
|
| 552 |
+
.b8 113
|
| 553 |
+
.b8 113
|
| 554 |
+
.b8 101
|
| 555 |
+
.b8 111
|
| 556 |
+
.b8 101
|
| 557 |
+
.b8 52
|
| 558 |
+
.b8 105
|
| 559 |
+
.b8 99
|
| 560 |
+
.b8 106
|
| 561 |
+
.b8 112
|
| 562 |
+
.b8 109
|
| 563 |
+
.b8 121
|
| 564 |
+
.b8 98
|
| 565 |
+
.b8 100
|
| 566 |
+
.b8 119
|
| 567 |
+
.b8 99
|
| 568 |
+
.b8 117
|
| 569 |
+
.b8 106
|
| 570 |
+
.b8 120
|
| 571 |
+
.b8 110
|
| 572 |
+
.b8 54
|
| 573 |
+
.b8 46
|
| 574 |
+
.b8 112
|
| 575 |
+
.b8 121
|
| 576 |
+
.b8 0
|
| 577 |
+
.b32 .debug_line // DW_AT_stmt_list
|
| 578 |
+
.b8 47 // DW_AT_comp_dir
|
| 579 |
+
.b8 119
|
| 580 |
+
.b8 111
|
| 581 |
+
.b8 114
|
| 582 |
+
.b8 107
|
| 583 |
+
.b8 115
|
| 584 |
+
.b8 112
|
| 585 |
+
.b8 97
|
| 586 |
+
.b8 99
|
| 587 |
+
.b8 101
|
| 588 |
+
.b8 47
|
| 589 |
+
.b8 104
|
| 590 |
+
.b8 97
|
| 591 |
+
.b8 110
|
| 592 |
+
.b8 114
|
| 593 |
+
.b8 117
|
| 594 |
+
.b8 105
|
| 595 |
+
.b8 47
|
| 596 |
+
.b8 83
|
| 597 |
+
.b8 112
|
| 598 |
+
.b8 101
|
| 599 |
+
.b8 99
|
| 600 |
+
.b8 70
|
| 601 |
+
.b8 111
|
| 602 |
+
.b8 114
|
| 603 |
+
.b8 103
|
| 604 |
+
.b8 101
|
| 605 |
+
.b8 45
|
| 606 |
+
.b8 101
|
| 607 |
+
.b8 120
|
| 608 |
+
.b8 116
|
| 609 |
+
.b8 47
|
| 610 |
+
.b8 99
|
| 611 |
+
.b8 97
|
| 612 |
+
.b8 99
|
| 613 |
+
.b8 104
|
| 614 |
+
.b8 101
|
| 615 |
+
.b8 47
|
| 616 |
+
.b8 99
|
| 617 |
+
.b8 111
|
| 618 |
+
.b8 109
|
| 619 |
+
.b8 112
|
| 620 |
+
.b8 105
|
| 621 |
+
.b8 108
|
| 622 |
+
.b8 101
|
| 623 |
+
.b8 100
|
| 624 |
+
.b8 95
|
| 625 |
+
.b8 107
|
| 626 |
+
.b8 101
|
| 627 |
+
.b8 114
|
| 628 |
+
.b8 110
|
| 629 |
+
.b8 101
|
| 630 |
+
.b8 108
|
| 631 |
+
.b8 115
|
| 632 |
+
.b8 47
|
| 633 |
+
.b8 53
|
| 634 |
+
.b8 97
|
| 635 |
+
.b8 0
|
| 636 |
+
.b8 2 // Abbrev [2] 0x8b:0x26 DW_TAG_subprogram
|
| 637 |
+
.b8 116 // DW_AT_name
|
| 638 |
+
.b8 114
|
| 639 |
+
.b8 105
|
| 640 |
+
.b8 116
|
| 641 |
+
.b8 111
|
| 642 |
+
.b8 110
|
| 643 |
+
.b8 95
|
| 644 |
+
.b8 114
|
| 645 |
+
.b8 101
|
| 646 |
+
.b8 100
|
| 647 |
+
.b8 95
|
| 648 |
+
.b8 102
|
| 649 |
+
.b8 117
|
| 650 |
+
.b8 115
|
| 651 |
+
.b8 101
|
| 652 |
+
.b8 100
|
| 653 |
+
.b8 95
|
| 654 |
+
.b8 95
|
| 655 |
+
.b8 116
|
| 656 |
+
.b8 111
|
| 657 |
+
.b8 95
|
| 658 |
+
.b8 99
|
| 659 |
+
.b8 111
|
| 660 |
+
.b8 112
|
| 661 |
+
.b8 121
|
| 662 |
+
.b8 95
|
| 663 |
+
.b8 109
|
| 664 |
+
.b8 117
|
| 665 |
+
.b8 108
|
| 666 |
+
.b8 95
|
| 667 |
+
.b8 115
|
| 668 |
+
.b8 117
|
| 669 |
+
.b8 109
|
| 670 |
+
.b8 95
|
| 671 |
+
.b8 48
|
| 672 |
+
.b8 0
|
| 673 |
+
.b8 1 // DW_AT_inline
|
| 674 |
+
.b8 3 // Abbrev [3] 0xb1:0x2e DW_TAG_subprogram
|
| 675 |
+
.b64 $L__func_begin0 // DW_AT_low_pc
|
| 676 |
+
.b64 $L__func_end0 // DW_AT_high_pc
|
| 677 |
+
.b32 139 // DW_AT_abstract_origin
|
| 678 |
+
.b8 4 // Abbrev [4] 0xc6:0x18 DW_TAG_inlined_subroutine
|
| 679 |
+
.b32 139 // DW_AT_abstract_origin
|
| 680 |
+
.b64 $L__tmp1 // DW_AT_low_pc
|
| 681 |
+
.b64 $L__tmp4 // DW_AT_high_pc
|
| 682 |
+
.b8 1 // DW_AT_call_file
|
| 683 |
+
.b8 51 // DW_AT_call_line
|
| 684 |
+
.b8 27 // DW_AT_call_column
|
| 685 |
+
.b8 0 // End Of Children Mark
|
| 686 |
+
.b8 0 // End Of Children Mark
|
| 687 |
+
}
|
| 688 |
+
.section .debug_macinfo { }
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.source
ADDED
|
@@ -0,0 +1,338 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":18:0)
|
| 2 |
+
#loc74 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":285:0)
|
| 3 |
+
#loc76 = loc(unknown)
|
| 4 |
+
#loc79 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":260:0)
|
| 5 |
+
#loc83 = loc("in_ptr0"(#loc))
|
| 6 |
+
#loc84 = loc("in_ptr1"(#loc))
|
| 7 |
+
#loc85 = loc("in_ptr2"(#loc))
|
| 8 |
+
#loc86 = loc("out_ptr0"(#loc))
|
| 9 |
+
#loc87 = loc("ks0"(#loc))
|
| 10 |
+
#loc88 = loc("ks1"(#loc))
|
| 11 |
+
#loc89 = loc("ks2"(#loc))
|
| 12 |
+
#loc90 = loc("xnumel"(#loc))
|
| 13 |
+
#loc91 = loc("r0_numel"(#loc))
|
| 14 |
+
#loc161 = loc("input"(#loc74))
|
| 15 |
+
#loc162 = loc("a"(#loc79))
|
| 16 |
+
#loc163 = loc("b"(#loc79))
|
| 17 |
+
module {
|
| 18 |
+
tt.func public @triton_red_fused__to_copy_mul_sum_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)), %in_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("in_ptr2"(#loc)), %out_ptr0: !tt.ptr<f32> {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 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
|
| 19 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc92)
|
| 20 |
+
%xoffset_0 = arith.constant 64 : i32 loc(#loc93)
|
| 21 |
+
%xoffset_1 = arith.constant 64 : i32 loc(#loc93)
|
| 22 |
+
%xoffset_2 = arith.muli %xoffset, %xoffset_1 : i32 loc(#loc93)
|
| 23 |
+
%xindex = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32> loc(#loc94)
|
| 24 |
+
%xindex_3 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<64xi32> -> tensor<64x1xi32> loc(#loc95)
|
| 25 |
+
%xindex_4 = tt.splat %xoffset_2 : i32 -> tensor<64x1xi32> loc(#loc96)
|
| 26 |
+
%xindex_5 = arith.addi %xindex_4, %xindex_3 : tensor<64x1xi32> loc(#loc96)
|
| 27 |
+
%xmask = tt.splat %xnumel : i32 -> tensor<64x1xi32> loc(#loc97)
|
| 28 |
+
%xmask_6 = arith.cmpi slt, %xindex_5, %xmask : tensor<64x1xi32> loc(#loc97)
|
| 29 |
+
%r0_base = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc98)
|
| 30 |
+
%r0_base_7 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<8xi32> -> tensor<1x8xi32> loc(#loc99)
|
| 31 |
+
%x1 = arith.extsi %xindex_5 : tensor<64x1xi32> to tensor<64x1xi64> loc(#loc100)
|
| 32 |
+
%x1_8 = tt.splat %ks0 : i64 -> tensor<64x1xi64> loc(#loc100)
|
| 33 |
+
%x1_9 = arith.divsi %x1, %x1_8 : tensor<64x1xi64> loc(#loc100)
|
| 34 |
+
%x0 = arith.extsi %xindex_5 : tensor<64x1xi32> to tensor<64x1xi64> loc(#loc101)
|
| 35 |
+
%x0_10 = tt.splat %ks0 : i64 -> tensor<64x1xi64> loc(#loc101)
|
| 36 |
+
%x0_11 = arith.remsi %x0, %x0_10 : tensor<64x1xi64> loc(#loc101)
|
| 37 |
+
%_tmp13 = arith.constant 0.000000e+00 : f32 loc(#loc102)
|
| 38 |
+
%_tmp13_12 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc102)
|
| 39 |
+
%c0_i32 = arith.constant 0 : i32 loc(#loc12)
|
| 40 |
+
%c8_i32 = arith.constant 8 : i32 loc(#loc12)
|
| 41 |
+
%0 = arith.bitcast %c0_i32 : i32 to i32 loc(#loc12)
|
| 42 |
+
%1 = arith.bitcast %r0_numel : i32 to i32 loc(#loc12)
|
| 43 |
+
%2 = arith.bitcast %c8_i32 : i32 to i32 loc(#loc12)
|
| 44 |
+
%3 = ub.poison : i32 loc(#loc12)
|
| 45 |
+
%_tmp13_13 = scf.for %r0_offset = %0 to %1 step %2 iter_args(%_tmp13_15 = %_tmp13_12) -> (tensor<64x8xf32>) : i32 {
|
| 46 |
+
%r0_index = tt.splat %r0_offset : i32 -> tensor<1x8xi32> loc(#loc104)
|
| 47 |
+
%r0_index_16 = arith.addi %r0_index, %r0_base_7 : tensor<1x8xi32> loc(#loc104)
|
| 48 |
+
%r0_mask = tt.splat %r0_numel : i32 -> tensor<1x8xi32> loc(#loc105)
|
| 49 |
+
%r0_mask_17 = arith.cmpi slt, %r0_index_16, %r0_mask : tensor<1x8xi32> loc(#loc105)
|
| 50 |
+
%tmp0 = arith.muli %ks1, %ks2 : i64 loc(#loc106)
|
| 51 |
+
%tmp0_18 = arith.constant 31 : i32 loc(#loc107)
|
| 52 |
+
%tmp0_19 = arith.constant 31 : i64 loc(#loc107)
|
| 53 |
+
%tmp0_20 = arith.addi %tmp0_19, %tmp0 : i64 loc(#loc107)
|
| 54 |
+
%tmp0_21 = arith.constant 32 : i32 loc(#loc108)
|
| 55 |
+
%tmp0_22 = arith.constant 32 : i64 loc(#loc108)
|
| 56 |
+
%tmp0_23 = arith.divsi %tmp0_20, %tmp0_22 : i64 loc(#loc108)
|
| 57 |
+
%tmp0_24 = tt.splat %tmp0_23 : i64 -> tensor<64x1xi64> loc(#loc109)
|
| 58 |
+
%tmp0_25 = arith.muli %x1_9, %tmp0_24 : tensor<64x1xi64> loc(#loc109)
|
| 59 |
+
%tmp0_26 = arith.extsi %r0_index_16 : tensor<1x8xi32> to tensor<1x8xi64> loc(#loc110)
|
| 60 |
+
%tmp0_27 = tt.broadcast %tmp0_26 : tensor<1x8xi64> -> tensor<64x8xi64> loc(#loc110)
|
| 61 |
+
%tmp0_28 = tt.broadcast %tmp0_25 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc110)
|
| 62 |
+
%tmp0_29 = arith.addi %tmp0_27, %tmp0_28 : tensor<64x8xi64> loc(#loc110)
|
| 63 |
+
%tmp1 = arith.muli %ks1, %ks2 : i64 loc(#loc111)
|
| 64 |
+
%tmp2 = tt.splat %tmp1 : i64 -> tensor<64x8xi64> loc(#loc112)
|
| 65 |
+
%tmp2_30 = arith.cmpi slt, %tmp0_29, %tmp2 : tensor<64x8xi64> loc(#loc112)
|
| 66 |
+
%tmp3 = arith.muli %ks1, %ks2 : i64 loc(#loc113)
|
| 67 |
+
%tmp3_31 = arith.constant 31 : i32 loc(#loc114)
|
| 68 |
+
%tmp3_32 = arith.constant 31 : i64 loc(#loc114)
|
| 69 |
+
%tmp3_33 = arith.addi %tmp3_32, %tmp3 : i64 loc(#loc114)
|
| 70 |
+
%tmp3_34 = arith.constant 32 : i32 loc(#loc115)
|
| 71 |
+
%tmp3_35 = arith.constant 32 : i64 loc(#loc115)
|
| 72 |
+
%tmp3_36 = arith.divsi %tmp3_33, %tmp3_35 : i64 loc(#loc115)
|
| 73 |
+
%tmp3_37 = tt.splat %tmp3_36 : i64 -> tensor<64x1xi64> loc(#loc116)
|
| 74 |
+
%tmp3_38 = arith.muli %x1_9, %tmp3_37 : tensor<64x1xi64> loc(#loc116)
|
| 75 |
+
%tmp3_39 = arith.extsi %r0_index_16 : tensor<1x8xi32> to tensor<1x8xi64> loc(#loc117)
|
| 76 |
+
%tmp3_40 = tt.broadcast %tmp3_39 : tensor<1x8xi64> -> tensor<64x8xi64> loc(#loc117)
|
| 77 |
+
%tmp3_41 = tt.broadcast %tmp3_38 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc117)
|
| 78 |
+
%tmp3_42 = arith.addi %tmp3_40, %tmp3_41 : tensor<64x8xi64> loc(#loc117)
|
| 79 |
+
%tmp3_43 = arith.muli %ks1, %ks2 : i64 loc(#loc118)
|
| 80 |
+
%tmp3_44 = tt.splat %tmp3_43 : i64 -> tensor<64x8xi64> loc(#loc119)
|
| 81 |
+
%tmp3_45 = arith.remsi %tmp3_42, %tmp3_44 : tensor<64x8xi64> loc(#loc119)
|
| 82 |
+
%tmp3_46 = tt.splat %ks0 : i64 -> tensor<64x8xi64> loc(#loc120)
|
| 83 |
+
%tmp3_47 = arith.muli %tmp3_46, %tmp3_45 : tensor<64x8xi64> loc(#loc120)
|
| 84 |
+
%tmp3_48 = tt.broadcast %x0_11 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc121)
|
| 85 |
+
%tmp3_49 = arith.addi %tmp3_48, %tmp3_47 : tensor<64x8xi64> loc(#loc121)
|
| 86 |
+
%tmp3_50 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc122)
|
| 87 |
+
%tmp3_51 = tt.addptr %tmp3_50, %tmp3_49 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi64> loc(#loc122)
|
| 88 |
+
%tmp3_52 = tt.broadcast %r0_mask_17 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc123)
|
| 89 |
+
%tmp3_53 = arith.andi %tmp3_52, %tmp2_30 : tensor<64x8xi1> loc(#loc123)
|
| 90 |
+
%tmp3_54 = tt.broadcast %xmask_6 : tensor<64x1xi1> -> tensor<64x8xi1> loc(#loc124)
|
| 91 |
+
%tmp3_55 = arith.andi %tmp3_53, %tmp3_54 : tensor<64x8xi1> loc(#loc124)
|
| 92 |
+
%tmp3_56 = arith.constant 0.000000e+00 : f32 loc(#loc125)
|
| 93 |
+
%tmp3_57 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc125)
|
| 94 |
+
%tmp3_58 = arith.truncf %tmp3_57 : tensor<64x8xf32> to tensor<64x8xbf16> loc(#loc125)
|
| 95 |
+
%tmp3_59 = tt.load %tmp3_51, %tmp3_55, %tmp3_58 evictionPolicy = evict_last : tensor<64x8x!tt.ptr<bf16>> loc(#loc125)
|
| 96 |
+
%tmp3_60 = arith.extf %tmp3_59 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc126)
|
| 97 |
+
%tmp4 = arith.muli %ks1, %ks2 : i64 loc(#loc127)
|
| 98 |
+
%tmp4_61 = arith.constant 31 : i32 loc(#loc128)
|
| 99 |
+
%tmp4_62 = arith.constant 31 : i64 loc(#loc128)
|
| 100 |
+
%tmp4_63 = arith.addi %tmp4_62, %tmp4 : i64 loc(#loc128)
|
| 101 |
+
%tmp4_64 = arith.constant 32 : i32 loc(#loc129)
|
| 102 |
+
%tmp4_65 = arith.constant 32 : i64 loc(#loc129)
|
| 103 |
+
%tmp4_66 = arith.divsi %tmp4_63, %tmp4_65 : i64 loc(#loc129)
|
| 104 |
+
%tmp4_67 = tt.splat %tmp4_66 : i64 -> tensor<64x1xi64> loc(#loc130)
|
| 105 |
+
%tmp4_68 = arith.muli %x1_9, %tmp4_67 : tensor<64x1xi64> loc(#loc130)
|
| 106 |
+
%tmp4_69 = arith.extsi %r0_index_16 : tensor<1x8xi32> to tensor<1x8xi64> loc(#loc131)
|
| 107 |
+
%tmp4_70 = tt.broadcast %tmp4_69 : tensor<1x8xi64> -> tensor<64x8xi64> loc(#loc131)
|
| 108 |
+
%tmp4_71 = tt.broadcast %tmp4_68 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc131)
|
| 109 |
+
%tmp4_72 = arith.addi %tmp4_70, %tmp4_71 : tensor<64x8xi64> loc(#loc131)
|
| 110 |
+
%tmp4_73 = arith.muli %ks1, %ks2 : i64 loc(#loc132)
|
| 111 |
+
%tmp4_74 = tt.splat %tmp4_73 : i64 -> tensor<64x8xi64> loc(#loc133)
|
| 112 |
+
%tmp4_75 = arith.remsi %tmp4_72, %tmp4_74 : tensor<64x8xi64> loc(#loc133)
|
| 113 |
+
%tmp4_76 = tt.splat %ks0 : i64 -> tensor<64x8xi64> loc(#loc134)
|
| 114 |
+
%tmp4_77 = arith.muli %tmp4_76, %tmp4_75 : tensor<64x8xi64> loc(#loc134)
|
| 115 |
+
%tmp4_78 = tt.broadcast %x0_11 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc135)
|
| 116 |
+
%tmp4_79 = arith.addi %tmp4_78, %tmp4_77 : tensor<64x8xi64> loc(#loc135)
|
| 117 |
+
%tmp4_80 = tt.splat %in_ptr1 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc136)
|
| 118 |
+
%tmp4_81 = tt.addptr %tmp4_80, %tmp4_79 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi64> loc(#loc136)
|
| 119 |
+
%tmp4_82 = tt.broadcast %r0_mask_17 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc137)
|
| 120 |
+
%tmp4_83 = arith.andi %tmp4_82, %tmp2_30 : tensor<64x8xi1> loc(#loc137)
|
| 121 |
+
%tmp4_84 = tt.broadcast %xmask_6 : tensor<64x1xi1> -> tensor<64x8xi1> loc(#loc138)
|
| 122 |
+
%tmp4_85 = arith.andi %tmp4_83, %tmp4_84 : tensor<64x8xi1> loc(#loc138)
|
| 123 |
+
%tmp4_86 = arith.constant 0.000000e+00 : f32 loc(#loc139)
|
| 124 |
+
%tmp4_87 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc139)
|
| 125 |
+
%tmp4_88 = arith.truncf %tmp4_87 : tensor<64x8xf32> to tensor<64x8xbf16> loc(#loc139)
|
| 126 |
+
%tmp4_89 = tt.load %tmp4_81, %tmp4_85, %tmp4_88 evictionPolicy = evict_last : tensor<64x8x!tt.ptr<bf16>> loc(#loc139)
|
| 127 |
+
%tmp4_90 = arith.extf %tmp4_89 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc140)
|
| 128 |
+
%tmp6 = arith.muli %ks1, %ks2 : i64 loc(#loc141)
|
| 129 |
+
%tmp6_91 = arith.constant 31 : i32 loc(#loc142)
|
| 130 |
+
%tmp6_92 = arith.constant 31 : i64 loc(#loc142)
|
| 131 |
+
%tmp6_93 = arith.addi %tmp6_92, %tmp6 : i64 loc(#loc142)
|
| 132 |
+
%tmp6_94 = arith.constant 32 : i32 loc(#loc143)
|
| 133 |
+
%tmp6_95 = arith.constant 32 : i64 loc(#loc143)
|
| 134 |
+
%tmp6_96 = arith.divsi %tmp6_93, %tmp6_95 : i64 loc(#loc143)
|
| 135 |
+
%tmp6_97 = tt.splat %tmp6_96 : i64 -> tensor<64x1xi64> loc(#loc144)
|
| 136 |
+
%tmp6_98 = arith.muli %x1_9, %tmp6_97 : tensor<64x1xi64> loc(#loc144)
|
| 137 |
+
%tmp6_99 = arith.extsi %r0_index_16 : tensor<1x8xi32> to tensor<1x8xi64> loc(#loc145)
|
| 138 |
+
%tmp6_100 = tt.broadcast %tmp6_99 : tensor<1x8xi64> -> tensor<64x8xi64> loc(#loc145)
|
| 139 |
+
%tmp6_101 = tt.broadcast %tmp6_98 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc145)
|
| 140 |
+
%tmp6_102 = arith.addi %tmp6_100, %tmp6_101 : tensor<64x8xi64> loc(#loc145)
|
| 141 |
+
%tmp6_103 = arith.muli %ks1, %ks2 : i64 loc(#loc146)
|
| 142 |
+
%tmp6_104 = tt.splat %tmp6_103 : i64 -> tensor<64x8xi64> loc(#loc147)
|
| 143 |
+
%tmp6_105 = arith.remsi %tmp6_102, %tmp6_104 : tensor<64x8xi64> loc(#loc147)
|
| 144 |
+
%tmp6_106 = tt.splat %in_ptr2 : !tt.ptr<f32> -> tensor<64x8x!tt.ptr<f32>> loc(#loc148)
|
| 145 |
+
%tmp6_107 = tt.addptr %tmp6_106, %tmp6_105 : tensor<64x8x!tt.ptr<f32>>, tensor<64x8xi64> loc(#loc148)
|
| 146 |
+
%tmp6_108 = tt.broadcast %r0_mask_17 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc149)
|
| 147 |
+
%tmp6_109 = arith.andi %tmp6_108, %tmp2_30 : tensor<64x8xi1> loc(#loc149)
|
| 148 |
+
%tmp6_110 = tt.broadcast %xmask_6 : tensor<64x1xi1> -> tensor<64x8xi1> loc(#loc150)
|
| 149 |
+
%tmp6_111 = arith.andi %tmp6_109, %tmp6_110 : tensor<64x8xi1> loc(#loc150)
|
| 150 |
+
%tmp6_112 = arith.constant 0.000000e+00 : f32 loc(#loc151)
|
| 151 |
+
%tmp6_113 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc151)
|
| 152 |
+
%tmp6_114 = tt.load %tmp6_107, %tmp6_111, %tmp6_113 evictionPolicy = evict_last : tensor<64x8x!tt.ptr<f32>> loc(#loc151)
|
| 153 |
+
%tmp7 = arith.mulf %tmp4_90, %tmp6_114 : tensor<64x8xf32> loc(#loc152)
|
| 154 |
+
%tmp9 = arith.mulf %tmp3_60, %tmp7 : tensor<64x8xf32> loc(#loc153)
|
| 155 |
+
%tmp10 = arith.constant 0.000000e+00 : f32 loc(#loc154)
|
| 156 |
+
%tmp10_115 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc154)
|
| 157 |
+
%tmp11 = arith.select %tmp2_30, %tmp9, %tmp10_115 : tensor<64x8xi1>, tensor<64x8xf32> loc(#loc155)
|
| 158 |
+
%tmp14 = arith.addf %_tmp13_15, %tmp11 : tensor<64x8xf32> loc(#loc156)
|
| 159 |
+
%_tmp13_116 = tt.broadcast %r0_mask_17 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc157)
|
| 160 |
+
%_tmp13_117 = tt.broadcast %xmask_6 : tensor<64x1xi1> -> tensor<64x8xi1> loc(#loc157)
|
| 161 |
+
%_tmp13_118 = arith.andi %_tmp13_116, %_tmp13_117 : tensor<64x8xi1> loc(#loc157)
|
| 162 |
+
%_tmp13_119 = arith.select %_tmp13_118, %tmp14, %_tmp13_15 : tensor<64x8xi1>, tensor<64x8xf32> loc(#loc158)
|
| 163 |
+
scf.yield %_tmp13_119 : tensor<64x8xf32> loc(#loc68)
|
| 164 |
+
} loc(#loc103)
|
| 165 |
+
%tmp13 = tt.call @"triton.language.standard.sum__fp32S64_8S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%_tmp13_13) : (tensor<64x8xf32>) -> tensor<64xf32> loc(#loc159)
|
| 166 |
+
%tmp13_14 = tt.expand_dims %tmp13 {axis = 1 : i32} : tensor<64xf32> -> tensor<64x1xf32> loc(#loc160)
|
| 167 |
+
%4 = tt.splat %out_ptr0 : !tt.ptr<f32> -> tensor<64x1x!tt.ptr<f32>> loc(#loc71)
|
| 168 |
+
%5 = tt.addptr %4, %xindex_5 : tensor<64x1x!tt.ptr<f32>>, tensor<64x1xi32> loc(#loc71)
|
| 169 |
+
tt.store %5, %tmp13_14, %xmask_6 : tensor<64x1x!tt.ptr<f32>> loc(#loc72)
|
| 170 |
+
tt.return loc(#loc73)
|
| 171 |
+
} loc(#loc)
|
| 172 |
+
tt.func private @"triton.language.standard.sum__fp32S64_8S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%input: tensor<64x8xf32> loc("input"(#loc74))) -> tensor<64xf32> attributes {noinline = false} {
|
| 173 |
+
%0 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
|
| 174 |
+
^bb0(%arg1: f32 loc(unknown), %arg2: f32 loc(unknown)):
|
| 175 |
+
%2 = tt.call @triton.language.standard._sum_combine__fp32_fp32__(%arg1, %arg2) : (f32, f32) -> f32 loc(#loc75)
|
| 176 |
+
tt.reduce.return %2 : f32 loc(#loc75)
|
| 177 |
+
}) : (tensor<64x8xf32>) -> tensor<64xf32> loc(#loc75)
|
| 178 |
+
tt.return %0 : tensor<64xf32> loc(#loc77)
|
| 179 |
+
^bb1: // no predecessors
|
| 180 |
+
%1 = ub.poison : tensor<64xf32> loc(#loc78)
|
| 181 |
+
tt.return %1 : tensor<64xf32> loc(#loc78)
|
| 182 |
+
} loc(#loc74)
|
| 183 |
+
tt.func private @triton.language.standard._sum_combine__fp32_fp32__(%a: f32 loc("a"(#loc79)), %b: f32 loc("b"(#loc79))) -> f32 attributes {noinline = false} {
|
| 184 |
+
%0 = arith.addf %a, %b : f32 loc(#loc80)
|
| 185 |
+
tt.return %0 : f32 loc(#loc81)
|
| 186 |
+
^bb1: // no predecessors
|
| 187 |
+
%1 = ub.poison : f32 loc(#loc82)
|
| 188 |
+
tt.return %1 : f32 loc(#loc82)
|
| 189 |
+
} loc(#loc79)
|
| 190 |
+
} loc(#loc)
|
| 191 |
+
#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":21:28)
|
| 192 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":21:33)
|
| 193 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:36)
|
| 194 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:44)
|
| 195 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:23)
|
| 196 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":23:21)
|
| 197 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":24:27)
|
| 198 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":24:37)
|
| 199 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":26:19)
|
| 200 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":27:19)
|
| 201 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":28:44)
|
| 202 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":30:40)
|
| 203 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":31:31)
|
| 204 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":32:29)
|
| 205 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:36)
|
| 206 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:32)
|
| 207 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:44)
|
| 208 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:26)
|
| 209 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:22)
|
| 210 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":37:19)
|
| 211 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":38:22)
|
| 212 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:67)
|
| 213 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:63)
|
| 214 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:75)
|
| 215 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:57)
|
| 216 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:53)
|
| 217 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:87)
|
| 218 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:83)
|
| 219 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:45)
|
| 220 |
+
#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:39)
|
| 221 |
+
#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:34)
|
| 222 |
+
#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:106)
|
| 223 |
+
#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:113)
|
| 224 |
+
#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:96)
|
| 225 |
+
#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:164)
|
| 226 |
+
#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:67)
|
| 227 |
+
#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:63)
|
| 228 |
+
#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:75)
|
| 229 |
+
#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:57)
|
| 230 |
+
#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:53)
|
| 231 |
+
#loc41 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:87)
|
| 232 |
+
#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:83)
|
| 233 |
+
#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:45)
|
| 234 |
+
#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:39)
|
| 235 |
+
#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:34)
|
| 236 |
+
#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:106)
|
| 237 |
+
#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:113)
|
| 238 |
+
#loc48 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:96)
|
| 239 |
+
#loc49 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:164)
|
| 240 |
+
#loc50 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:57)
|
| 241 |
+
#loc51 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:53)
|
| 242 |
+
#loc52 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:65)
|
| 243 |
+
#loc53 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:47)
|
| 244 |
+
#loc54 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:43)
|
| 245 |
+
#loc55 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:77)
|
| 246 |
+
#loc56 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:73)
|
| 247 |
+
#loc57 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:35)
|
| 248 |
+
#loc58 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:95)
|
| 249 |
+
#loc59 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:102)
|
| 250 |
+
#loc60 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:85)
|
| 251 |
+
#loc61 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":43:22)
|
| 252 |
+
#loc62 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":45:22)
|
| 253 |
+
#loc63 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":46:39)
|
| 254 |
+
#loc64 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":47:37)
|
| 255 |
+
#loc65 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":49:25)
|
| 256 |
+
#loc66 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:36)
|
| 257 |
+
#loc67 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:50)
|
| 258 |
+
#loc68 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:8)
|
| 259 |
+
#loc69 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":51:27)
|
| 260 |
+
#loc70 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":51:30)
|
| 261 |
+
#loc71 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:25)
|
| 262 |
+
#loc72 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:37)
|
| 263 |
+
#loc73 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:4)
|
| 264 |
+
#loc75 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 265 |
+
#loc77 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:11)
|
| 266 |
+
#loc78 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:4)
|
| 267 |
+
#loc80 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 268 |
+
#loc81 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:11)
|
| 269 |
+
#loc82 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:4)
|
| 270 |
+
#loc92 = loc("xoffset"(#loc1))
|
| 271 |
+
#loc93 = loc("xoffset"(#loc2))
|
| 272 |
+
#loc94 = loc("xindex"(#loc3))
|
| 273 |
+
#loc95 = loc("xindex"(#loc4))
|
| 274 |
+
#loc96 = loc("xindex"(#loc5))
|
| 275 |
+
#loc97 = loc("xmask"(#loc6))
|
| 276 |
+
#loc98 = loc("r0_base"(#loc7))
|
| 277 |
+
#loc99 = loc("r0_base"(#loc8))
|
| 278 |
+
#loc100 = loc("x1"(#loc9))
|
| 279 |
+
#loc101 = loc("x0"(#loc10))
|
| 280 |
+
#loc102 = loc("_tmp13"(#loc11))
|
| 281 |
+
#loc103 = loc("_tmp13"(#loc12))
|
| 282 |
+
#loc104 = loc("r0_index"(#loc13))
|
| 283 |
+
#loc105 = loc("r0_mask"(#loc14))
|
| 284 |
+
#loc106 = loc("tmp0"(#loc15))
|
| 285 |
+
#loc107 = loc("tmp0"(#loc16))
|
| 286 |
+
#loc108 = loc("tmp0"(#loc17))
|
| 287 |
+
#loc109 = loc("tmp0"(#loc18))
|
| 288 |
+
#loc110 = loc("tmp0"(#loc19))
|
| 289 |
+
#loc111 = loc("tmp1"(#loc20))
|
| 290 |
+
#loc112 = loc("tmp2"(#loc21))
|
| 291 |
+
#loc113 = loc("tmp3"(#loc22))
|
| 292 |
+
#loc114 = loc("tmp3"(#loc23))
|
| 293 |
+
#loc115 = loc("tmp3"(#loc24))
|
| 294 |
+
#loc116 = loc("tmp3"(#loc25))
|
| 295 |
+
#loc117 = loc("tmp3"(#loc26))
|
| 296 |
+
#loc118 = loc("tmp3"(#loc27))
|
| 297 |
+
#loc119 = loc("tmp3"(#loc28))
|
| 298 |
+
#loc120 = loc("tmp3"(#loc29))
|
| 299 |
+
#loc121 = loc("tmp3"(#loc30))
|
| 300 |
+
#loc122 = loc("tmp3"(#loc31))
|
| 301 |
+
#loc123 = loc("tmp3"(#loc32))
|
| 302 |
+
#loc124 = loc("tmp3"(#loc33))
|
| 303 |
+
#loc125 = loc("tmp3"(#loc34))
|
| 304 |
+
#loc126 = loc("tmp3"(#loc35))
|
| 305 |
+
#loc127 = loc("tmp4"(#loc36))
|
| 306 |
+
#loc128 = loc("tmp4"(#loc37))
|
| 307 |
+
#loc129 = loc("tmp4"(#loc38))
|
| 308 |
+
#loc130 = loc("tmp4"(#loc39))
|
| 309 |
+
#loc131 = loc("tmp4"(#loc40))
|
| 310 |
+
#loc132 = loc("tmp4"(#loc41))
|
| 311 |
+
#loc133 = loc("tmp4"(#loc42))
|
| 312 |
+
#loc134 = loc("tmp4"(#loc43))
|
| 313 |
+
#loc135 = loc("tmp4"(#loc44))
|
| 314 |
+
#loc136 = loc("tmp4"(#loc45))
|
| 315 |
+
#loc137 = loc("tmp4"(#loc46))
|
| 316 |
+
#loc138 = loc("tmp4"(#loc47))
|
| 317 |
+
#loc139 = loc("tmp4"(#loc48))
|
| 318 |
+
#loc140 = loc("tmp4"(#loc49))
|
| 319 |
+
#loc141 = loc("tmp6"(#loc50))
|
| 320 |
+
#loc142 = loc("tmp6"(#loc51))
|
| 321 |
+
#loc143 = loc("tmp6"(#loc52))
|
| 322 |
+
#loc144 = loc("tmp6"(#loc53))
|
| 323 |
+
#loc145 = loc("tmp6"(#loc54))
|
| 324 |
+
#loc146 = loc("tmp6"(#loc55))
|
| 325 |
+
#loc147 = loc("tmp6"(#loc56))
|
| 326 |
+
#loc148 = loc("tmp6"(#loc57))
|
| 327 |
+
#loc149 = loc("tmp6"(#loc58))
|
| 328 |
+
#loc150 = loc("tmp6"(#loc59))
|
| 329 |
+
#loc151 = loc("tmp6"(#loc60))
|
| 330 |
+
#loc152 = loc("tmp7"(#loc61))
|
| 331 |
+
#loc153 = loc("tmp9"(#loc62))
|
| 332 |
+
#loc154 = loc("tmp10"(#loc63))
|
| 333 |
+
#loc155 = loc("tmp11"(#loc64))
|
| 334 |
+
#loc156 = loc("tmp14"(#loc65))
|
| 335 |
+
#loc157 = loc("_tmp13"(#loc66))
|
| 336 |
+
#loc158 = loc("_tmp13"(#loc67))
|
| 337 |
+
#loc159 = loc("tmp13"(#loc69))
|
| 338 |
+
#loc160 = loc("tmp13"(#loc70))
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ttgir
ADDED
|
@@ -0,0 +1,176 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#blocked = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [2, 2], order = [0, 1]}>
|
| 2 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":18:0)
|
| 3 |
+
#loc1 = loc(unknown)
|
| 4 |
+
#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":51:27)
|
| 5 |
+
#loc46 = loc("in_ptr0"(#loc))
|
| 6 |
+
#loc47 = loc("in_ptr1"(#loc))
|
| 7 |
+
#loc48 = loc("in_ptr2"(#loc))
|
| 8 |
+
#loc49 = loc("out_ptr0"(#loc))
|
| 9 |
+
#loc50 = loc("ks0"(#loc))
|
| 10 |
+
#loc51 = loc("ks1"(#loc))
|
| 11 |
+
#loc52 = loc("ks2"(#loc))
|
| 12 |
+
#loc53 = loc("xnumel"(#loc))
|
| 13 |
+
#loc54 = loc("r0_numel"(#loc))
|
| 14 |
+
#loc91 = loc("tmp13"(#loc40))
|
| 15 |
+
#loc94 = loc(callsite(#loc1 at #loc91))
|
| 16 |
+
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
|
| 17 |
+
tt.func public @triton_red_fused__to_copy_mul_sum_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)), %in_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("in_ptr2"(#loc)), %out_ptr0: !tt.ptr<f32> {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 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
|
| 18 |
+
%cst = arith.constant dense<0.000000e+00> : tensor<64x8xf32, #blocked> loc(#loc1)
|
| 19 |
+
%c31_i64 = arith.constant 31 : i64 loc(#loc1)
|
| 20 |
+
%c32_i64 = arith.constant 32 : i64 loc(#loc1)
|
| 21 |
+
%c0_i32 = arith.constant 0 : i32 loc(#loc1)
|
| 22 |
+
%c8_i32 = arith.constant 8 : i32 loc(#loc1)
|
| 23 |
+
%c64_i32 = arith.constant 64 : i32 loc(#loc1)
|
| 24 |
+
%cst_0 = arith.constant dense<0.000000e+00> : tensor<64x8xbf16, #blocked> loc(#loc1)
|
| 25 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc55)
|
| 26 |
+
%xoffset_1 = arith.muli %xoffset, %c64_i32 : i32 loc(#loc56)
|
| 27 |
+
%xindex = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc57)
|
| 28 |
+
%xindex_2 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<64x1xi32, #blocked> loc(#loc57)
|
| 29 |
+
%xindex_3 = tt.splat %xoffset_1 : i32 -> tensor<64x1xi32, #blocked> loc(#loc58)
|
| 30 |
+
%xindex_4 = arith.addi %xindex_3, %xindex_2 : tensor<64x1xi32, #blocked> loc(#loc58)
|
| 31 |
+
%xmask = tt.splat %xnumel : i32 -> tensor<64x1xi32, #blocked> loc(#loc59)
|
| 32 |
+
%xmask_5 = arith.cmpi slt, %xindex_4, %xmask : tensor<64x1xi32, #blocked> loc(#loc59)
|
| 33 |
+
%r0_base = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc60)
|
| 34 |
+
%r0_base_6 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<8xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x8xi32, #blocked> loc(#loc60)
|
| 35 |
+
%x1 = arith.extsi %xindex_4 : tensor<64x1xi32, #blocked> to tensor<64x1xi64, #blocked> loc(#loc61)
|
| 36 |
+
%x1_7 = tt.splat %ks0 : i64 -> tensor<64x1xi64, #blocked> loc(#loc61)
|
| 37 |
+
%x1_8 = arith.divsi %x1, %x1_7 : tensor<64x1xi64, #blocked> loc(#loc61)
|
| 38 |
+
%x0 = arith.remsi %x1, %x1_7 : tensor<64x1xi64, #blocked> loc(#loc62)
|
| 39 |
+
%r0_mask = tt.splat %r0_numel : i32 -> tensor<1x8xi32, #blocked> loc(#loc63)
|
| 40 |
+
%tmp0 = arith.muli %ks1, %ks2 : i64 loc(#loc64)
|
| 41 |
+
%tmp0_9 = arith.addi %tmp0, %c31_i64 : i64 loc(#loc65)
|
| 42 |
+
%tmp0_10 = arith.divsi %tmp0_9, %c32_i64 : i64 loc(#loc66)
|
| 43 |
+
%tmp0_11 = tt.splat %tmp0_10 : i64 -> tensor<64x1xi64, #blocked> loc(#loc67)
|
| 44 |
+
%tmp0_12 = arith.muli %x1_8, %tmp0_11 : tensor<64x1xi64, #blocked> loc(#loc67)
|
| 45 |
+
%tmp0_13 = tt.broadcast %tmp0_12 : tensor<64x1xi64, #blocked> -> tensor<64x8xi64, #blocked> loc(#loc68)
|
| 46 |
+
%tmp2 = tt.splat %tmp0 : i64 -> tensor<64x8xi64, #blocked> loc(#loc69)
|
| 47 |
+
%tmp3 = tt.splat %ks0 : i64 -> tensor<64x8xi64, #blocked> loc(#loc70)
|
| 48 |
+
%tmp3_14 = tt.broadcast %x0 : tensor<64x1xi64, #blocked> -> tensor<64x8xi64, #blocked> loc(#loc71)
|
| 49 |
+
%tmp3_15 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc72)
|
| 50 |
+
%tmp3_16 = tt.broadcast %xmask_5 : tensor<64x1xi1, #blocked> -> tensor<64x8xi1, #blocked> loc(#loc73)
|
| 51 |
+
%tmp4 = tt.splat %in_ptr1 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc74)
|
| 52 |
+
%tmp6 = tt.splat %in_ptr2 : !tt.ptr<f32> -> tensor<64x8x!tt.ptr<f32>, #blocked> loc(#loc75)
|
| 53 |
+
%_tmp13 = scf.for %_tmp13_18 = %c0_i32 to %r0_numel step %c8_i32 iter_args(%arg10 = %cst) -> (tensor<64x8xf32, #blocked>) : i32 {
|
| 54 |
+
%r0_index = tt.splat %_tmp13_18 : i32 -> tensor<1x8xi32, #blocked> loc(#loc77)
|
| 55 |
+
%r0_index_19 = arith.addi %r0_index, %r0_base_6 : tensor<1x8xi32, #blocked> loc(#loc77)
|
| 56 |
+
%r0_mask_20 = arith.cmpi slt, %r0_index_19, %r0_mask : tensor<1x8xi32, #blocked> loc(#loc63)
|
| 57 |
+
%tmp0_21 = arith.extsi %r0_index_19 : tensor<1x8xi32, #blocked> to tensor<1x8xi64, #blocked> loc(#loc68)
|
| 58 |
+
%tmp0_22 = tt.broadcast %tmp0_21 : tensor<1x8xi64, #blocked> -> tensor<64x8xi64, #blocked> loc(#loc68)
|
| 59 |
+
%tmp0_23 = arith.addi %tmp0_22, %tmp0_13 : tensor<64x8xi64, #blocked> loc(#loc68)
|
| 60 |
+
%tmp2_24 = arith.cmpi slt, %tmp0_23, %tmp2 : tensor<64x8xi64, #blocked> loc(#loc69)
|
| 61 |
+
%tmp3_25 = arith.remsi %tmp0_23, %tmp2 : tensor<64x8xi64, #blocked> loc(#loc78)
|
| 62 |
+
%tmp3_26 = arith.muli %tmp3, %tmp3_25 : tensor<64x8xi64, #blocked> loc(#loc70)
|
| 63 |
+
%tmp3_27 = arith.addi %tmp3_14, %tmp3_26 : tensor<64x8xi64, #blocked> loc(#loc71)
|
| 64 |
+
%tmp3_28 = tt.addptr %tmp3_15, %tmp3_27 : tensor<64x8x!tt.ptr<bf16>, #blocked>, tensor<64x8xi64, #blocked> loc(#loc72)
|
| 65 |
+
%tmp3_29 = tt.broadcast %r0_mask_20 : tensor<1x8xi1, #blocked> -> tensor<64x8xi1, #blocked> loc(#loc79)
|
| 66 |
+
%tmp3_30 = arith.andi %tmp3_29, %tmp2_24 : tensor<64x8xi1, #blocked> loc(#loc79)
|
| 67 |
+
%tmp3_31 = arith.andi %tmp3_30, %tmp3_16 : tensor<64x8xi1, #blocked> loc(#loc73)
|
| 68 |
+
%tmp3_32 = tt.load %tmp3_28, %tmp3_31, %cst_0 evictionPolicy = evict_last : tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc80)
|
| 69 |
+
%tmp3_33 = arith.extf %tmp3_32 : tensor<64x8xbf16, #blocked> to tensor<64x8xf32, #blocked> loc(#loc81)
|
| 70 |
+
%tmp4_34 = tt.addptr %tmp4, %tmp3_27 : tensor<64x8x!tt.ptr<bf16>, #blocked>, tensor<64x8xi64, #blocked> loc(#loc74)
|
| 71 |
+
%tmp4_35 = tt.load %tmp4_34, %tmp3_31, %cst_0 evictionPolicy = evict_last : tensor<64x8x!tt.ptr<bf16>, #blocked> loc(#loc82)
|
| 72 |
+
%tmp4_36 = arith.extf %tmp4_35 : tensor<64x8xbf16, #blocked> to tensor<64x8xf32, #blocked> loc(#loc83)
|
| 73 |
+
%tmp6_37 = tt.addptr %tmp6, %tmp3_25 : tensor<64x8x!tt.ptr<f32>, #blocked>, tensor<64x8xi64, #blocked> loc(#loc75)
|
| 74 |
+
%tmp6_38 = tt.load %tmp6_37, %tmp3_31, %cst evictionPolicy = evict_last : tensor<64x8x!tt.ptr<f32>, #blocked> loc(#loc84)
|
| 75 |
+
%tmp7 = arith.mulf %tmp4_36, %tmp6_38 : tensor<64x8xf32, #blocked> loc(#loc85)
|
| 76 |
+
%tmp9 = arith.mulf %tmp3_33, %tmp7 : tensor<64x8xf32, #blocked> loc(#loc86)
|
| 77 |
+
%tmp11 = arith.select %tmp2_24, %tmp9, %cst : tensor<64x8xi1, #blocked>, tensor<64x8xf32, #blocked> loc(#loc87)
|
| 78 |
+
%tmp14 = arith.addf %arg10, %tmp11 : tensor<64x8xf32, #blocked> loc(#loc88)
|
| 79 |
+
%_tmp13_39 = arith.andi %tmp3_29, %tmp3_16 : tensor<64x8xi1, #blocked> loc(#loc89)
|
| 80 |
+
%_tmp13_40 = arith.select %_tmp13_39, %tmp14, %arg10 : tensor<64x8xi1, #blocked>, tensor<64x8xf32, #blocked> loc(#loc90)
|
| 81 |
+
scf.yield %_tmp13_40 : tensor<64x8xf32, #blocked> loc(#loc38)
|
| 82 |
+
} loc(#loc76)
|
| 83 |
+
%tmp13 = "tt.reduce"(%_tmp13) <{axis = 1 : i32}> ({
|
| 84 |
+
^bb0(%tmp13_18: f32 loc(callsite(#loc1 at #loc91)), %tmp13_19: f32 loc(callsite(#loc1 at #loc91))):
|
| 85 |
+
%tmp13_20 = arith.addf %tmp13_18, %tmp13_19 : f32 loc(#loc95)
|
| 86 |
+
tt.reduce.return %tmp13_20 : f32 loc(#loc93)
|
| 87 |
+
}) : (tensor<64x8xf32, #blocked>) -> tensor<64xf32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc93)
|
| 88 |
+
%tmp13_17 = tt.expand_dims %tmp13 {axis = 1 : i32} : tensor<64xf32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<64x1xf32, #blocked> loc(#loc92)
|
| 89 |
+
%0 = tt.splat %out_ptr0 : !tt.ptr<f32> -> tensor<64x1x!tt.ptr<f32>, #blocked> loc(#loc43)
|
| 90 |
+
%1 = tt.addptr %0, %xindex_4 : tensor<64x1x!tt.ptr<f32>, #blocked>, tensor<64x1xi32, #blocked> loc(#loc43)
|
| 91 |
+
tt.store %1, %tmp13_17, %xmask_5 : tensor<64x1x!tt.ptr<f32>, #blocked> loc(#loc44)
|
| 92 |
+
tt.return loc(#loc45)
|
| 93 |
+
} loc(#loc)
|
| 94 |
+
} loc(#loc)
|
| 95 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":21:28)
|
| 96 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":21:33)
|
| 97 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:44)
|
| 98 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:23)
|
| 99 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":23:21)
|
| 100 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":24:37)
|
| 101 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":26:19)
|
| 102 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":27:19)
|
| 103 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":32:29)
|
| 104 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:36)
|
| 105 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:32)
|
| 106 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:44)
|
| 107 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:26)
|
| 108 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:22)
|
| 109 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":38:22)
|
| 110 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:45)
|
| 111 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:39)
|
| 112 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:34)
|
| 113 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:113)
|
| 114 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:34)
|
| 115 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:35)
|
| 116 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":30:40)
|
| 117 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":31:31)
|
| 118 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:83)
|
| 119 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:106)
|
| 120 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:96)
|
| 121 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:164)
|
| 122 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:96)
|
| 123 |
+
#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:164)
|
| 124 |
+
#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:85)
|
| 125 |
+
#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":43:22)
|
| 126 |
+
#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":45:22)
|
| 127 |
+
#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":47:37)
|
| 128 |
+
#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":49:25)
|
| 129 |
+
#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:36)
|
| 130 |
+
#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:50)
|
| 131 |
+
#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:8)
|
| 132 |
+
#loc39 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 133 |
+
#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 134 |
+
#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":51:30)
|
| 135 |
+
#loc43 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:25)
|
| 136 |
+
#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:37)
|
| 137 |
+
#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:4)
|
| 138 |
+
#loc55 = loc("xoffset"(#loc2))
|
| 139 |
+
#loc56 = loc("xoffset"(#loc3))
|
| 140 |
+
#loc57 = loc("xindex"(#loc4))
|
| 141 |
+
#loc58 = loc("xindex"(#loc5))
|
| 142 |
+
#loc59 = loc("xmask"(#loc6))
|
| 143 |
+
#loc60 = loc("r0_base"(#loc7))
|
| 144 |
+
#loc61 = loc("x1"(#loc8))
|
| 145 |
+
#loc62 = loc("x0"(#loc9))
|
| 146 |
+
#loc63 = loc("r0_mask"(#loc10))
|
| 147 |
+
#loc64 = loc("tmp0"(#loc11))
|
| 148 |
+
#loc65 = loc("tmp0"(#loc12))
|
| 149 |
+
#loc66 = loc("tmp0"(#loc13))
|
| 150 |
+
#loc67 = loc("tmp0"(#loc14))
|
| 151 |
+
#loc68 = loc("tmp0"(#loc15))
|
| 152 |
+
#loc69 = loc("tmp2"(#loc16))
|
| 153 |
+
#loc70 = loc("tmp3"(#loc17))
|
| 154 |
+
#loc71 = loc("tmp3"(#loc18))
|
| 155 |
+
#loc72 = loc("tmp3"(#loc19))
|
| 156 |
+
#loc73 = loc("tmp3"(#loc20))
|
| 157 |
+
#loc74 = loc("tmp4"(#loc21))
|
| 158 |
+
#loc75 = loc("tmp6"(#loc22))
|
| 159 |
+
#loc76 = loc("_tmp13"(#loc23))
|
| 160 |
+
#loc77 = loc("r0_index"(#loc24))
|
| 161 |
+
#loc78 = loc("tmp3"(#loc25))
|
| 162 |
+
#loc79 = loc("tmp3"(#loc26))
|
| 163 |
+
#loc80 = loc("tmp3"(#loc27))
|
| 164 |
+
#loc81 = loc("tmp3"(#loc28))
|
| 165 |
+
#loc82 = loc("tmp4"(#loc29))
|
| 166 |
+
#loc83 = loc("tmp4"(#loc30))
|
| 167 |
+
#loc84 = loc("tmp6"(#loc31))
|
| 168 |
+
#loc85 = loc("tmp7"(#loc32))
|
| 169 |
+
#loc86 = loc("tmp9"(#loc33))
|
| 170 |
+
#loc87 = loc("tmp11"(#loc34))
|
| 171 |
+
#loc88 = loc("tmp14"(#loc35))
|
| 172 |
+
#loc89 = loc("_tmp13"(#loc36))
|
| 173 |
+
#loc90 = loc("_tmp13"(#loc37))
|
| 174 |
+
#loc92 = loc("tmp13"(#loc42))
|
| 175 |
+
#loc93 = loc(callsite(#loc39 at #loc91))
|
| 176 |
+
#loc95 = loc(callsite(#loc41 at #loc93))
|
SpecForge-ext/cache/compiled_kernels/triton/1/5KYA3IXELFPSJGJX4ZJXZEGZVFALDHEANQ7YQX554R55QVAC6F5A/triton_red_fused__to_copy_mul_sum_0.ttir
ADDED
|
@@ -0,0 +1,179 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":18:0)
|
| 2 |
+
#loc1 = loc(unknown)
|
| 3 |
+
#loc42 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":51:27)
|
| 4 |
+
#loc48 = loc("in_ptr0"(#loc))
|
| 5 |
+
#loc49 = loc("in_ptr1"(#loc))
|
| 6 |
+
#loc50 = loc("in_ptr2"(#loc))
|
| 7 |
+
#loc51 = loc("out_ptr0"(#loc))
|
| 8 |
+
#loc52 = loc("ks0"(#loc))
|
| 9 |
+
#loc53 = loc("ks1"(#loc))
|
| 10 |
+
#loc54 = loc("ks2"(#loc))
|
| 11 |
+
#loc55 = loc("xnumel"(#loc))
|
| 12 |
+
#loc56 = loc("r0_numel"(#loc))
|
| 13 |
+
#loc95 = loc("tmp13"(#loc42))
|
| 14 |
+
#loc98 = loc(callsite(#loc1 at #loc95))
|
| 15 |
+
module {
|
| 16 |
+
tt.func public @triton_red_fused__to_copy_mul_sum_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)), %in_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("in_ptr2"(#loc)), %out_ptr0: !tt.ptr<f32> {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 {tt.divisibility = 16 : i32} loc("xnumel"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
|
| 17 |
+
%cst = arith.constant dense<0.000000e+00> : tensor<64x8xbf16> loc(#loc1)
|
| 18 |
+
%c8_i32 = arith.constant 8 : i32 loc(#loc2)
|
| 19 |
+
%c0_i32 = arith.constant 0 : i32 loc(#loc2)
|
| 20 |
+
%c32_i64 = arith.constant 32 : i64 loc(#loc1)
|
| 21 |
+
%c31_i64 = arith.constant 31 : i64 loc(#loc1)
|
| 22 |
+
%cst_0 = arith.constant dense<0.000000e+00> : tensor<64x8xf32> loc(#loc1)
|
| 23 |
+
%c64_i32 = arith.constant 64 : i32 loc(#loc1)
|
| 24 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc57)
|
| 25 |
+
%xoffset_1 = arith.muli %xoffset, %c64_i32 : i32 loc(#loc58)
|
| 26 |
+
%xindex = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32> loc(#loc59)
|
| 27 |
+
%xindex_2 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<64xi32> -> tensor<64x1xi32> loc(#loc60)
|
| 28 |
+
%xindex_3 = tt.splat %xoffset_1 : i32 -> tensor<64x1xi32> loc(#loc61)
|
| 29 |
+
%xindex_4 = arith.addi %xindex_3, %xindex_2 : tensor<64x1xi32> loc(#loc61)
|
| 30 |
+
%xmask = tt.splat %xnumel : i32 -> tensor<64x1xi32> loc(#loc62)
|
| 31 |
+
%xmask_5 = arith.cmpi slt, %xindex_4, %xmask : tensor<64x1xi32> loc(#loc62)
|
| 32 |
+
%r0_base = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc63)
|
| 33 |
+
%r0_base_6 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<8xi32> -> tensor<1x8xi32> loc(#loc64)
|
| 34 |
+
%x1 = arith.extsi %xindex_4 : tensor<64x1xi32> to tensor<64x1xi64> loc(#loc65)
|
| 35 |
+
%x1_7 = tt.splat %ks0 : i64 -> tensor<64x1xi64> loc(#loc65)
|
| 36 |
+
%x1_8 = arith.divsi %x1, %x1_7 : tensor<64x1xi64> loc(#loc65)
|
| 37 |
+
%x0 = arith.remsi %x1, %x1_7 : tensor<64x1xi64> loc(#loc66)
|
| 38 |
+
%_tmp13 = scf.for %r0_offset = %c0_i32 to %r0_numel step %c8_i32 iter_args(%_tmp13_10 = %cst_0) -> (tensor<64x8xf32>) : i32 {
|
| 39 |
+
%r0_index = tt.splat %r0_offset : i32 -> tensor<1x8xi32> loc(#loc68)
|
| 40 |
+
%r0_index_11 = arith.addi %r0_index, %r0_base_6 : tensor<1x8xi32> loc(#loc68)
|
| 41 |
+
%r0_mask = tt.splat %r0_numel : i32 -> tensor<1x8xi32> loc(#loc69)
|
| 42 |
+
%r0_mask_12 = arith.cmpi slt, %r0_index_11, %r0_mask : tensor<1x8xi32> loc(#loc69)
|
| 43 |
+
%tmp0 = arith.muli %ks1, %ks2 : i64 loc(#loc70)
|
| 44 |
+
%tmp0_13 = arith.addi %tmp0, %c31_i64 : i64 loc(#loc71)
|
| 45 |
+
%tmp0_14 = arith.divsi %tmp0_13, %c32_i64 : i64 loc(#loc72)
|
| 46 |
+
%tmp0_15 = tt.splat %tmp0_14 : i64 -> tensor<64x1xi64> loc(#loc73)
|
| 47 |
+
%tmp0_16 = arith.muli %x1_8, %tmp0_15 : tensor<64x1xi64> loc(#loc73)
|
| 48 |
+
%tmp0_17 = arith.extsi %r0_index_11 : tensor<1x8xi32> to tensor<1x8xi64> loc(#loc74)
|
| 49 |
+
%tmp0_18 = tt.broadcast %tmp0_17 : tensor<1x8xi64> -> tensor<64x8xi64> loc(#loc74)
|
| 50 |
+
%tmp0_19 = tt.broadcast %tmp0_16 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc74)
|
| 51 |
+
%tmp0_20 = arith.addi %tmp0_18, %tmp0_19 : tensor<64x8xi64> loc(#loc74)
|
| 52 |
+
%tmp2 = tt.splat %tmp0 : i64 -> tensor<64x8xi64> loc(#loc75)
|
| 53 |
+
%tmp2_21 = arith.cmpi slt, %tmp0_20, %tmp2 : tensor<64x8xi64> loc(#loc75)
|
| 54 |
+
%tmp3 = arith.remsi %tmp0_20, %tmp2 : tensor<64x8xi64> loc(#loc76)
|
| 55 |
+
%tmp3_22 = tt.splat %ks0 : i64 -> tensor<64x8xi64> loc(#loc77)
|
| 56 |
+
%tmp3_23 = arith.muli %tmp3_22, %tmp3 : tensor<64x8xi64> loc(#loc77)
|
| 57 |
+
%tmp3_24 = tt.broadcast %x0 : tensor<64x1xi64> -> tensor<64x8xi64> loc(#loc78)
|
| 58 |
+
%tmp3_25 = arith.addi %tmp3_24, %tmp3_23 : tensor<64x8xi64> loc(#loc78)
|
| 59 |
+
%tmp3_26 = tt.splat %in_ptr0 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc79)
|
| 60 |
+
%tmp3_27 = tt.addptr %tmp3_26, %tmp3_25 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi64> loc(#loc79)
|
| 61 |
+
%tmp3_28 = tt.broadcast %r0_mask_12 : tensor<1x8xi1> -> tensor<64x8xi1> loc(#loc80)
|
| 62 |
+
%tmp3_29 = arith.andi %tmp3_28, %tmp2_21 : tensor<64x8xi1> loc(#loc80)
|
| 63 |
+
%tmp3_30 = tt.broadcast %xmask_5 : tensor<64x1xi1> -> tensor<64x8xi1> loc(#loc81)
|
| 64 |
+
%tmp3_31 = arith.andi %tmp3_29, %tmp3_30 : tensor<64x8xi1> loc(#loc81)
|
| 65 |
+
%tmp3_32 = tt.load %tmp3_27, %tmp3_31, %cst evictionPolicy = evict_last : tensor<64x8x!tt.ptr<bf16>> loc(#loc82)
|
| 66 |
+
%tmp3_33 = arith.extf %tmp3_32 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc83)
|
| 67 |
+
%tmp4 = tt.splat %in_ptr1 : !tt.ptr<bf16> -> tensor<64x8x!tt.ptr<bf16>> loc(#loc84)
|
| 68 |
+
%tmp4_34 = tt.addptr %tmp4, %tmp3_25 : tensor<64x8x!tt.ptr<bf16>>, tensor<64x8xi64> loc(#loc84)
|
| 69 |
+
%tmp4_35 = tt.load %tmp4_34, %tmp3_31, %cst evictionPolicy = evict_last : tensor<64x8x!tt.ptr<bf16>> loc(#loc85)
|
| 70 |
+
%tmp4_36 = arith.extf %tmp4_35 : tensor<64x8xbf16> to tensor<64x8xf32> loc(#loc86)
|
| 71 |
+
%tmp6 = tt.splat %in_ptr2 : !tt.ptr<f32> -> tensor<64x8x!tt.ptr<f32>> loc(#loc87)
|
| 72 |
+
%tmp6_37 = tt.addptr %tmp6, %tmp3 : tensor<64x8x!tt.ptr<f32>>, tensor<64x8xi64> loc(#loc87)
|
| 73 |
+
%tmp6_38 = tt.load %tmp6_37, %tmp3_31, %cst_0 evictionPolicy = evict_last : tensor<64x8x!tt.ptr<f32>> loc(#loc88)
|
| 74 |
+
%tmp7 = arith.mulf %tmp4_36, %tmp6_38 : tensor<64x8xf32> loc(#loc89)
|
| 75 |
+
%tmp9 = arith.mulf %tmp3_33, %tmp7 : tensor<64x8xf32> loc(#loc90)
|
| 76 |
+
%tmp11 = arith.select %tmp2_21, %tmp9, %cst_0 : tensor<64x8xi1>, tensor<64x8xf32> loc(#loc91)
|
| 77 |
+
%tmp14 = arith.addf %_tmp13_10, %tmp11 : tensor<64x8xf32> loc(#loc92)
|
| 78 |
+
%_tmp13_39 = arith.andi %tmp3_28, %tmp3_30 : tensor<64x8xi1> loc(#loc93)
|
| 79 |
+
%_tmp13_40 = arith.select %_tmp13_39, %tmp14, %_tmp13_10 : tensor<64x8xi1>, tensor<64x8xf32> loc(#loc94)
|
| 80 |
+
scf.yield %_tmp13_40 : tensor<64x8xf32> loc(#loc40)
|
| 81 |
+
} loc(#loc67)
|
| 82 |
+
%tmp13 = "tt.reduce"(%_tmp13) <{axis = 1 : i32}> ({
|
| 83 |
+
^bb0(%tmp13_10: f32 loc(callsite(#loc1 at #loc95)), %tmp13_11: f32 loc(callsite(#loc1 at #loc95))):
|
| 84 |
+
%tmp13_12 = arith.addf %tmp13_10, %tmp13_11 : f32 loc(#loc99)
|
| 85 |
+
tt.reduce.return %tmp13_12 : f32 loc(#loc97)
|
| 86 |
+
}) : (tensor<64x8xf32>) -> tensor<64xf32> loc(#loc97)
|
| 87 |
+
%tmp13_9 = tt.expand_dims %tmp13 {axis = 1 : i32} : tensor<64xf32> -> tensor<64x1xf32> loc(#loc96)
|
| 88 |
+
%0 = tt.splat %out_ptr0 : !tt.ptr<f32> -> tensor<64x1x!tt.ptr<f32>> loc(#loc45)
|
| 89 |
+
%1 = tt.addptr %0, %xindex_4 : tensor<64x1x!tt.ptr<f32>>, tensor<64x1xi32> loc(#loc45)
|
| 90 |
+
tt.store %1, %tmp13_9, %xmask_5 : tensor<64x1x!tt.ptr<f32>> loc(#loc46)
|
| 91 |
+
tt.return loc(#loc47)
|
| 92 |
+
} loc(#loc)
|
| 93 |
+
} loc(#loc)
|
| 94 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":30:40)
|
| 95 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":21:28)
|
| 96 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":21:33)
|
| 97 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:36)
|
| 98 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:44)
|
| 99 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":22:23)
|
| 100 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":23:21)
|
| 101 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":24:27)
|
| 102 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":24:37)
|
| 103 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":26:19)
|
| 104 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":27:19)
|
| 105 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":31:31)
|
| 106 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":32:29)
|
| 107 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:36)
|
| 108 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:32)
|
| 109 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:44)
|
| 110 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:26)
|
| 111 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":36:22)
|
| 112 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":38:22)
|
| 113 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:83)
|
| 114 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:45)
|
| 115 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:39)
|
| 116 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:34)
|
| 117 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:106)
|
| 118 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:113)
|
| 119 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:96)
|
| 120 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":39:164)
|
| 121 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:34)
|
| 122 |
+
#loc30 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:96)
|
| 123 |
+
#loc31 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":40:164)
|
| 124 |
+
#loc32 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:35)
|
| 125 |
+
#loc33 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":42:85)
|
| 126 |
+
#loc34 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":43:22)
|
| 127 |
+
#loc35 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":45:22)
|
| 128 |
+
#loc36 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":47:37)
|
| 129 |
+
#loc37 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":49:25)
|
| 130 |
+
#loc38 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:36)
|
| 131 |
+
#loc39 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:50)
|
| 132 |
+
#loc40 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":50:8)
|
| 133 |
+
#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 134 |
+
#loc43 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 135 |
+
#loc44 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":51:30)
|
| 136 |
+
#loc45 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:25)
|
| 137 |
+
#loc46 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:37)
|
| 138 |
+
#loc47 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/5a/c5aj4zurklc4jhvqpnfghi3vz6khjwkqqeoe4icjpmybdwcujxn6.py":52:4)
|
| 139 |
+
#loc57 = loc("xoffset"(#loc3))
|
| 140 |
+
#loc58 = loc("xoffset"(#loc4))
|
| 141 |
+
#loc59 = loc("xindex"(#loc5))
|
| 142 |
+
#loc60 = loc("xindex"(#loc6))
|
| 143 |
+
#loc61 = loc("xindex"(#loc7))
|
| 144 |
+
#loc62 = loc("xmask"(#loc8))
|
| 145 |
+
#loc63 = loc("r0_base"(#loc9))
|
| 146 |
+
#loc64 = loc("r0_base"(#loc10))
|
| 147 |
+
#loc65 = loc("x1"(#loc11))
|
| 148 |
+
#loc66 = loc("x0"(#loc12))
|
| 149 |
+
#loc67 = loc("_tmp13"(#loc2))
|
| 150 |
+
#loc68 = loc("r0_index"(#loc13))
|
| 151 |
+
#loc69 = loc("r0_mask"(#loc14))
|
| 152 |
+
#loc70 = loc("tmp0"(#loc15))
|
| 153 |
+
#loc71 = loc("tmp0"(#loc16))
|
| 154 |
+
#loc72 = loc("tmp0"(#loc17))
|
| 155 |
+
#loc73 = loc("tmp0"(#loc18))
|
| 156 |
+
#loc74 = loc("tmp0"(#loc19))
|
| 157 |
+
#loc75 = loc("tmp2"(#loc20))
|
| 158 |
+
#loc76 = loc("tmp3"(#loc21))
|
| 159 |
+
#loc77 = loc("tmp3"(#loc22))
|
| 160 |
+
#loc78 = loc("tmp3"(#loc23))
|
| 161 |
+
#loc79 = loc("tmp3"(#loc24))
|
| 162 |
+
#loc80 = loc("tmp3"(#loc25))
|
| 163 |
+
#loc81 = loc("tmp3"(#loc26))
|
| 164 |
+
#loc82 = loc("tmp3"(#loc27))
|
| 165 |
+
#loc83 = loc("tmp3"(#loc28))
|
| 166 |
+
#loc84 = loc("tmp4"(#loc29))
|
| 167 |
+
#loc85 = loc("tmp4"(#loc30))
|
| 168 |
+
#loc86 = loc("tmp4"(#loc31))
|
| 169 |
+
#loc87 = loc("tmp6"(#loc32))
|
| 170 |
+
#loc88 = loc("tmp6"(#loc33))
|
| 171 |
+
#loc89 = loc("tmp7"(#loc34))
|
| 172 |
+
#loc90 = loc("tmp9"(#loc35))
|
| 173 |
+
#loc91 = loc("tmp11"(#loc36))
|
| 174 |
+
#loc92 = loc("tmp14"(#loc37))
|
| 175 |
+
#loc93 = loc("_tmp13"(#loc38))
|
| 176 |
+
#loc94 = loc("_tmp13"(#loc39))
|
| 177 |
+
#loc96 = loc("tmp13"(#loc44))
|
| 178 |
+
#loc97 = loc(callsite(#loc41 at #loc95))
|
| 179 |
+
#loc99 = loc(callsite(#loc43 at #loc97))
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/__grp__triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.source", "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttir", "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttgir", "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.llir", "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ptx", "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.cubin", "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.cubin
ADDED
|
Binary file (7.96 kB). View file
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"hash": "ec0409dee971ea6f361feab9e50436845d14fe0c95c533fcd5da5d2041e210f9", "target": {"backend": "cuda", "arch": 90, "warp_size": 32}, "num_warps": 2, "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_per_fused_clamp_min_div_eq_mul_squeeze_sum_4"}
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.llir
ADDED
|
@@ -0,0 +1,91 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
; ModuleID = 'LLVMDialectModule'
|
| 2 |
+
source_filename = "LLVMDialectModule"
|
| 3 |
+
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"
|
| 4 |
+
|
| 5 |
+
; Function Attrs: nounwind
|
| 6 |
+
define ptx_kernel void @triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, ptr addrspace(1) readnone captures(none) %4, ptr addrspace(1) readnone captures(none) %5) local_unnamed_addr #0 !dbg !4 {
|
| 7 |
+
%7 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !7
|
| 8 |
+
%8 = and i32 %7, 1, !dbg !7
|
| 9 |
+
%9 = zext nneg i32 %8 to i64, !dbg !8
|
| 10 |
+
%10 = getelementptr i64, ptr addrspace(1) %0, i64 %9, !dbg !8
|
| 11 |
+
%11 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09ld.global.b64 { $0 }, [ $1 + 0 ];", "=l,l"(ptr addrspace(1) %10) #4, !dbg !9
|
| 12 |
+
%12 = getelementptr i64, ptr addrspace(1) %1, i64 %9, !dbg !10
|
| 13 |
+
%13 = tail call i64 asm sideeffect "mov.u64 $0, 0x0;\0A\09ld.global.b64 { $0 }, [ $1 + 0 ];", "=l,l"(ptr addrspace(1) %12) #4, !dbg !11
|
| 14 |
+
%extelt.offset = lshr i64 %11, 32, !dbg !12
|
| 15 |
+
%14 = trunc nuw i64 %extelt.offset to i32, !dbg !12
|
| 16 |
+
%15 = trunc i64 %11 to i32, !dbg !12
|
| 17 |
+
%16 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %15, i32 1, i32 31), !dbg !12
|
| 18 |
+
%17 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %14, i32 1, i32 31), !dbg !12
|
| 19 |
+
%18 = insertelement <2 x i32> poison, i32 %16, i64 0, !dbg !12
|
| 20 |
+
%19 = insertelement <2 x i32> %18, i32 %17, i64 1, !dbg !12
|
| 21 |
+
%20 = bitcast <2 x i32> %19 to i64, !dbg !12
|
| 22 |
+
%21 = add i64 %11, %20, !dbg !16
|
| 23 |
+
%extelt.offset1 = lshr i64 %13, 32, !dbg !17
|
| 24 |
+
%22 = trunc nuw i64 %extelt.offset1 to i32, !dbg !17
|
| 25 |
+
%23 = trunc i64 %13 to i32, !dbg !17
|
| 26 |
+
%24 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %23, i32 1, i32 31), !dbg !17
|
| 27 |
+
%25 = tail call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %22, i32 1, i32 31), !dbg !17
|
| 28 |
+
%26 = insertelement <2 x i32> poison, i32 %24, i64 0, !dbg !17
|
| 29 |
+
%27 = insertelement <2 x i32> %26, i32 %25, i64 1, !dbg !17
|
| 30 |
+
%28 = bitcast <2 x i32> %27 to i64, !dbg !17
|
| 31 |
+
%29 = add i64 %13, %28, !dbg !19
|
| 32 |
+
%30 = sitofp i64 %21 to float, !dbg !20
|
| 33 |
+
%31 = sitofp i64 %29 to float, !dbg !21
|
| 34 |
+
%32 = icmp sgt i64 %29, 0, !dbg !22
|
| 35 |
+
%33 = select i1 %32, float %31, float 0x3EB0C6F7A0000000, !dbg !26
|
| 36 |
+
%34 = tail call float @llvm.nvvm.div.full(float %30, float %33), !dbg !27
|
| 37 |
+
%35 = and i32 %7, 63, !dbg !28
|
| 38 |
+
%36 = icmp eq i32 %35, 0, !dbg !28
|
| 39 |
+
%37 = bitcast float %34 to i32, !dbg !28
|
| 40 |
+
tail call void asm sideeffect "@$2 st.global.b32 [ $1 + 0 ], { $0 };", "r,l,b"(i32 %37, ptr addrspace(1) %2, i1 %36) #4, !dbg !28
|
| 41 |
+
ret void, !dbg !29
|
| 42 |
+
}
|
| 43 |
+
|
| 44 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 45 |
+
declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
|
| 46 |
+
|
| 47 |
+
; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
|
| 48 |
+
declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) #2
|
| 49 |
+
|
| 50 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(none)
|
| 51 |
+
declare float @llvm.nvvm.div.full(float, float) #3
|
| 52 |
+
|
| 53 |
+
attributes #0 = { nounwind "nvvm.reqntid"="64" }
|
| 54 |
+
attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
|
| 55 |
+
attributes #2 = { convergent nocallback nounwind memory(inaccessiblemem: readwrite) }
|
| 56 |
+
attributes #3 = { mustprogress nocallback nofree nosync nounwind willreturn memory(none) }
|
| 57 |
+
attributes #4 = { nounwind }
|
| 58 |
+
|
| 59 |
+
!llvm.dbg.cu = !{!0}
|
| 60 |
+
!llvm.module.flags = !{!2, !3}
|
| 61 |
+
|
| 62 |
+
!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
|
| 63 |
+
!1 = !DIFile(filename: "ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5")
|
| 64 |
+
!2 = !{i32 2, !"Debug Info Version", i32 3}
|
| 65 |
+
!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
|
| 66 |
+
!4 = distinct !DISubprogram(name: "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4", linkageName: "triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
|
| 67 |
+
!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
|
| 68 |
+
!6 = !{}
|
| 69 |
+
!7 = !DILocation(line: 27, column: 38, scope: !4)
|
| 70 |
+
!8 = !DILocation(line: 33, column: 30, scope: !4)
|
| 71 |
+
!9 = !DILocation(line: 33, column: 37, scope: !4)
|
| 72 |
+
!10 = !DILocation(line: 34, column: 30, scope: !4)
|
| 73 |
+
!11 = !DILocation(line: 34, column: 37, scope: !4)
|
| 74 |
+
!12 = !DILocation(line: 291, column: 36, scope: !13, inlinedAt: !15)
|
| 75 |
+
!13 = distinct !DILexicalBlockFile(scope: !4, file: !14, discriminator: 0)
|
| 76 |
+
!14 = !DIFile(filename: "standard.py", directory: "/workspace/specforge/lib/python3.11/site-packages/triton/language")
|
| 77 |
+
!15 = !DILocation(line: 36, column: 24, scope: !4)
|
| 78 |
+
!16 = !DILocation(line: 261, column: 15, scope: !13, inlinedAt: !15)
|
| 79 |
+
!17 = !DILocation(line: 291, column: 36, scope: !13, inlinedAt: !18)
|
| 80 |
+
!18 = !DILocation(line: 38, column: 24, scope: !4)
|
| 81 |
+
!19 = !DILocation(line: 261, column: 15, scope: !13, inlinedAt: !18)
|
| 82 |
+
!20 = !DILocation(line: 39, column: 19, scope: !4)
|
| 83 |
+
!21 = !DILocation(line: 40, column: 19, scope: !4)
|
| 84 |
+
!22 = !DILocation(line: 110, column: 15, scope: !23, inlinedAt: !25)
|
| 85 |
+
!23 = distinct !DILexicalBlockFile(scope: !4, file: !24, discriminator: 0)
|
| 86 |
+
!24 = !DIFile(filename: "triton_helpers.py", directory: "/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime")
|
| 87 |
+
!25 = !DILocation(line: 42, column: 41, scope: !4)
|
| 88 |
+
!26 = !DILocation(line: 113, column: 29, scope: !23, inlinedAt: !25)
|
| 89 |
+
!27 = !DILocation(line: 43, column: 20, scope: !4)
|
| 90 |
+
!28 = !DILocation(line: 44, column: 68, scope: !4)
|
| 91 |
+
!29 = !DILocation(line: 44, column: 4, scope: !4)
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ptx
ADDED
|
@@ -0,0 +1,373 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
//
|
| 2 |
+
// Generated by LLVM NVPTX Back-End
|
| 3 |
+
//
|
| 4 |
+
|
| 5 |
+
.version 8.7
|
| 6 |
+
.target sm_90a
|
| 7 |
+
.address_size 64
|
| 8 |
+
|
| 9 |
+
// .globl triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4 // -- Begin function triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4
|
| 10 |
+
// @triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4
|
| 11 |
+
.visible .entry triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4(
|
| 12 |
+
.param .u64 .ptr .global .align 1 triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_0,
|
| 13 |
+
.param .u64 .ptr .global .align 1 triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_1,
|
| 14 |
+
.param .u64 .ptr .global .align 1 triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_2,
|
| 15 |
+
.param .u32 triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_3,
|
| 16 |
+
.param .u64 .ptr .global .align 1 triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_4,
|
| 17 |
+
.param .u64 .ptr .global .align 1 triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_5
|
| 18 |
+
)
|
| 19 |
+
.reqntid 64
|
| 20 |
+
{
|
| 21 |
+
.reg .pred %p<3>;
|
| 22 |
+
.reg .b32 %r<16>;
|
| 23 |
+
.reg .b64 %rd<19>;
|
| 24 |
+
.loc 1 18 0 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:18:0
|
| 25 |
+
$L__func_begin0:
|
| 26 |
+
.loc 1 18 0 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:18:0
|
| 27 |
+
|
| 28 |
+
// %bb.0:
|
| 29 |
+
ld.param.b64 %rd6, [triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_0];
|
| 30 |
+
ld.param.b64 %rd7, [triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_1];
|
| 31 |
+
$L__tmp0:
|
| 32 |
+
.loc 1 27 38 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:27:38
|
| 33 |
+
mov.u32 %r2, %tid.x;
|
| 34 |
+
and.b32 %r3, %r2, 1;
|
| 35 |
+
ld.param.b64 %rd5, [triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4_param_2];
|
| 36 |
+
.loc 1 33 30 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:33:30
|
| 37 |
+
mul.wide.u32 %rd8, %r3, 8;
|
| 38 |
+
add.s64 %rd2, %rd6, %rd8;
|
| 39 |
+
.loc 1 33 37 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:33:37
|
| 40 |
+
// begin inline asm
|
| 41 |
+
mov.u64 %rd1, 0x0;
|
| 42 |
+
ld.global.b64 { %rd1 }, [ %rd2 + 0 ];
|
| 43 |
+
// end inline asm
|
| 44 |
+
.loc 1 34 30 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:34:30
|
| 45 |
+
add.s64 %rd4, %rd7, %rd8;
|
| 46 |
+
.loc 1 34 37 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:34:37
|
| 47 |
+
// begin inline asm
|
| 48 |
+
mov.u64 %rd3, 0x0;
|
| 49 |
+
ld.global.b64 { %rd3 }, [ %rd4 + 0 ];
|
| 50 |
+
// end inline asm
|
| 51 |
+
$L__tmp1:
|
| 52 |
+
.loc 2 291 36 // standard.py:291:36 @[ ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:36:24 ]
|
| 53 |
+
mov.b64 {_, %r4}, %rd1;
|
| 54 |
+
cvt.u32.u64 %r5, %rd1;
|
| 55 |
+
shfl.sync.bfly.b32 %r6, %r5, 1, 31, -1;
|
| 56 |
+
shfl.sync.bfly.b32 %r7, %r4, 1, 31, -1;
|
| 57 |
+
cvt.u64.u32 %rd9, %r6;
|
| 58 |
+
cvt.u64.u32 %rd10, %r7;
|
| 59 |
+
shl.b64 %rd11, %rd10, 32;
|
| 60 |
+
or.b64 %rd12, %rd9, %rd11;
|
| 61 |
+
.loc 2 261 15 // standard.py:261:15 @[ ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:36:24 ]
|
| 62 |
+
add.s64 %rd13, %rd1, %rd12;
|
| 63 |
+
$L__tmp2:
|
| 64 |
+
.loc 2 291 36 // standard.py:291:36 @[ ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:38:24 ]
|
| 65 |
+
mov.b64 {_, %r8}, %rd3;
|
| 66 |
+
cvt.u32.u64 %r9, %rd3;
|
| 67 |
+
shfl.sync.bfly.b32 %r10, %r9, 1, 31, -1;
|
| 68 |
+
shfl.sync.bfly.b32 %r11, %r8, 1, 31, -1;
|
| 69 |
+
cvt.u64.u32 %rd14, %r10;
|
| 70 |
+
cvt.u64.u32 %rd15, %r11;
|
| 71 |
+
shl.b64 %rd16, %rd15, 32;
|
| 72 |
+
or.b64 %rd17, %rd14, %rd16;
|
| 73 |
+
.loc 2 261 15 // standard.py:261:15 @[ ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:38:24 ]
|
| 74 |
+
add.s64 %rd18, %rd3, %rd17;
|
| 75 |
+
$L__tmp3:
|
| 76 |
+
.loc 1 39 19 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:39:19
|
| 77 |
+
cvt.rn.f32.s64 %r12, %rd13;
|
| 78 |
+
.loc 1 40 19 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:40:19
|
| 79 |
+
cvt.rn.f32.s64 %r13, %rd18;
|
| 80 |
+
$L__tmp4:
|
| 81 |
+
.loc 3 110 15 // triton_helpers.py:110:15 @[ ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:42:41 ]
|
| 82 |
+
setp.gt.s64 %p2, %rd18, 0;
|
| 83 |
+
.loc 3 113 29 // triton_helpers.py:113:29 @[ ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:42:41 ]
|
| 84 |
+
selp.f32 %r14, %r13, 0f358637BD, %p2;
|
| 85 |
+
$L__tmp5:
|
| 86 |
+
.loc 1 43 20 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:43:20
|
| 87 |
+
div.full.f32 %r1, %r12, %r14;
|
| 88 |
+
.loc 1 44 68 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:44:68
|
| 89 |
+
and.b32 %r15, %r2, 63;
|
| 90 |
+
setp.eq.b32 %p1, %r15, 0;
|
| 91 |
+
// begin inline asm
|
| 92 |
+
@%p1 st.global.b32 [ %rd5 + 0 ], { %r1 };
|
| 93 |
+
// end inline asm
|
| 94 |
+
.loc 1 44 4 // ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py:44:4
|
| 95 |
+
ret;
|
| 96 |
+
$L__tmp6:
|
| 97 |
+
$L__func_end0:
|
| 98 |
+
// -- End function
|
| 99 |
+
}
|
| 100 |
+
.file 1 "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py"
|
| 101 |
+
.file 2 "/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py"
|
| 102 |
+
.file 3 "/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py"
|
| 103 |
+
.section .debug_abbrev
|
| 104 |
+
{
|
| 105 |
+
.b8 1 // Abbreviation Code
|
| 106 |
+
.b8 17 // DW_TAG_compile_unit
|
| 107 |
+
.b8 1 // DW_CHILDREN_yes
|
| 108 |
+
.b8 37 // DW_AT_producer
|
| 109 |
+
.b8 8 // DW_FORM_string
|
| 110 |
+
.b8 19 // DW_AT_language
|
| 111 |
+
.b8 5 // DW_FORM_data2
|
| 112 |
+
.b8 3 // DW_AT_name
|
| 113 |
+
.b8 8 // DW_FORM_string
|
| 114 |
+
.b8 16 // DW_AT_stmt_list
|
| 115 |
+
.b8 6 // DW_FORM_data4
|
| 116 |
+
.b8 27 // DW_AT_comp_dir
|
| 117 |
+
.b8 8 // DW_FORM_string
|
| 118 |
+
.b8 0 // EOM(1)
|
| 119 |
+
.b8 0 // EOM(2)
|
| 120 |
+
.b8 2 // Abbreviation Code
|
| 121 |
+
.b8 46 // DW_TAG_subprogram
|
| 122 |
+
.b8 0 // DW_CHILDREN_no
|
| 123 |
+
.b8 3 // DW_AT_name
|
| 124 |
+
.b8 8 // DW_FORM_string
|
| 125 |
+
.b8 32 // DW_AT_inline
|
| 126 |
+
.b8 11 // DW_FORM_data1
|
| 127 |
+
.b8 0 // EOM(1)
|
| 128 |
+
.b8 0 // EOM(2)
|
| 129 |
+
.b8 3 // Abbreviation Code
|
| 130 |
+
.b8 46 // DW_TAG_subprogram
|
| 131 |
+
.b8 1 // DW_CHILDREN_yes
|
| 132 |
+
.b8 17 // DW_AT_low_pc
|
| 133 |
+
.b8 1 // DW_FORM_addr
|
| 134 |
+
.b8 18 // DW_AT_high_pc
|
| 135 |
+
.b8 1 // DW_FORM_addr
|
| 136 |
+
.b8 49 // DW_AT_abstract_origin
|
| 137 |
+
.b8 19 // DW_FORM_ref4
|
| 138 |
+
.b8 0 // EOM(1)
|
| 139 |
+
.b8 0 // EOM(2)
|
| 140 |
+
.b8 4 // Abbreviation Code
|
| 141 |
+
.b8 29 // DW_TAG_inlined_subroutine
|
| 142 |
+
.b8 0 // DW_CHILDREN_no
|
| 143 |
+
.b8 49 // DW_AT_abstract_origin
|
| 144 |
+
.b8 19 // DW_FORM_ref4
|
| 145 |
+
.b8 17 // DW_AT_low_pc
|
| 146 |
+
.b8 1 // DW_FORM_addr
|
| 147 |
+
.b8 18 // DW_AT_high_pc
|
| 148 |
+
.b8 1 // DW_FORM_addr
|
| 149 |
+
.b8 88 // DW_AT_call_file
|
| 150 |
+
.b8 11 // DW_FORM_data1
|
| 151 |
+
.b8 89 // DW_AT_call_line
|
| 152 |
+
.b8 11 // DW_FORM_data1
|
| 153 |
+
.b8 87 // DW_AT_call_column
|
| 154 |
+
.b8 11 // DW_FORM_data1
|
| 155 |
+
.b8 0 // EOM(1)
|
| 156 |
+
.b8 0 // EOM(2)
|
| 157 |
+
.b8 0 // EOM(3)
|
| 158 |
+
}
|
| 159 |
+
.section .debug_info
|
| 160 |
+
{
|
| 161 |
+
.b32 284 // Length of Unit
|
| 162 |
+
.b8 2 // DWARF version number
|
| 163 |
+
.b8 0
|
| 164 |
+
.b32 .debug_abbrev // Offset Into Abbrev. Section
|
| 165 |
+
.b8 8 // Address Size (in bytes)
|
| 166 |
+
.b8 1 // Abbrev [1] 0xb:0x115 DW_TAG_compile_unit
|
| 167 |
+
.b8 116 // DW_AT_producer
|
| 168 |
+
.b8 114
|
| 169 |
+
.b8 105
|
| 170 |
+
.b8 116
|
| 171 |
+
.b8 111
|
| 172 |
+
.b8 110
|
| 173 |
+
.b8 0
|
| 174 |
+
.b8 2 // DW_AT_language
|
| 175 |
+
.b8 0
|
| 176 |
+
.b8 99 // DW_AT_name
|
| 177 |
+
.b8 104
|
| 178 |
+
.b8 53
|
| 179 |
+
.b8 113
|
| 180 |
+
.b8 116
|
| 181 |
+
.b8 115
|
| 182 |
+
.b8 102
|
| 183 |
+
.b8 101
|
| 184 |
+
.b8 121
|
| 185 |
+
.b8 122
|
| 186 |
+
.b8 111
|
| 187 |
+
.b8 112
|
| 188 |
+
.b8 121
|
| 189 |
+
.b8 103
|
| 190 |
+
.b8 102
|
| 191 |
+
.b8 110
|
| 192 |
+
.b8 100
|
| 193 |
+
.b8 114
|
| 194 |
+
.b8 102
|
| 195 |
+
.b8 53
|
| 196 |
+
.b8 52
|
| 197 |
+
.b8 104
|
| 198 |
+
.b8 111
|
| 199 |
+
.b8 99
|
| 200 |
+
.b8 102
|
| 201 |
+
.b8 54
|
| 202 |
+
.b8 114
|
| 203 |
+
.b8 100
|
| 204 |
+
.b8 103
|
| 205 |
+
.b8 55
|
| 206 |
+
.b8 50
|
| 207 |
+
.b8 112
|
| 208 |
+
.b8 100
|
| 209 |
+
.b8 106
|
| 210 |
+
.b8 52
|
| 211 |
+
.b8 101
|
| 212 |
+
.b8 112
|
| 213 |
+
.b8 122
|
| 214 |
+
.b8 117
|
| 215 |
+
.b8 101
|
| 216 |
+
.b8 97
|
| 217 |
+
.b8 103
|
| 218 |
+
.b8 102
|
| 219 |
+
.b8 53
|
| 220 |
+
.b8 118
|
| 221 |
+
.b8 122
|
| 222 |
+
.b8 53
|
| 223 |
+
.b8 122
|
| 224 |
+
.b8 115
|
| 225 |
+
.b8 121
|
| 226 |
+
.b8 98
|
| 227 |
+
.b8 55
|
| 228 |
+
.b8 46
|
| 229 |
+
.b8 112
|
| 230 |
+
.b8 121
|
| 231 |
+
.b8 0
|
| 232 |
+
.b32 .debug_line // DW_AT_stmt_list
|
| 233 |
+
.b8 47 // DW_AT_comp_dir
|
| 234 |
+
.b8 119
|
| 235 |
+
.b8 111
|
| 236 |
+
.b8 114
|
| 237 |
+
.b8 107
|
| 238 |
+
.b8 115
|
| 239 |
+
.b8 112
|
| 240 |
+
.b8 97
|
| 241 |
+
.b8 99
|
| 242 |
+
.b8 101
|
| 243 |
+
.b8 47
|
| 244 |
+
.b8 104
|
| 245 |
+
.b8 97
|
| 246 |
+
.b8 110
|
| 247 |
+
.b8 114
|
| 248 |
+
.b8 117
|
| 249 |
+
.b8 105
|
| 250 |
+
.b8 47
|
| 251 |
+
.b8 83
|
| 252 |
+
.b8 112
|
| 253 |
+
.b8 101
|
| 254 |
+
.b8 99
|
| 255 |
+
.b8 70
|
| 256 |
+
.b8 111
|
| 257 |
+
.b8 114
|
| 258 |
+
.b8 103
|
| 259 |
+
.b8 101
|
| 260 |
+
.b8 45
|
| 261 |
+
.b8 101
|
| 262 |
+
.b8 120
|
| 263 |
+
.b8 116
|
| 264 |
+
.b8 47
|
| 265 |
+
.b8 99
|
| 266 |
+
.b8 97
|
| 267 |
+
.b8 99
|
| 268 |
+
.b8 104
|
| 269 |
+
.b8 101
|
| 270 |
+
.b8 47
|
| 271 |
+
.b8 99
|
| 272 |
+
.b8 111
|
| 273 |
+
.b8 109
|
| 274 |
+
.b8 112
|
| 275 |
+
.b8 105
|
| 276 |
+
.b8 108
|
| 277 |
+
.b8 101
|
| 278 |
+
.b8 100
|
| 279 |
+
.b8 95
|
| 280 |
+
.b8 107
|
| 281 |
+
.b8 101
|
| 282 |
+
.b8 114
|
| 283 |
+
.b8 110
|
| 284 |
+
.b8 101
|
| 285 |
+
.b8 108
|
| 286 |
+
.b8 115
|
| 287 |
+
.b8 47
|
| 288 |
+
.b8 104
|
| 289 |
+
.b8 53
|
| 290 |
+
.b8 0
|
| 291 |
+
.b8 2 // Abbrev [2] 0x8b:0x36 DW_TAG_subprogram
|
| 292 |
+
.b8 116 // DW_AT_name
|
| 293 |
+
.b8 114
|
| 294 |
+
.b8 105
|
| 295 |
+
.b8 116
|
| 296 |
+
.b8 111
|
| 297 |
+
.b8 110
|
| 298 |
+
.b8 95
|
| 299 |
+
.b8 112
|
| 300 |
+
.b8 101
|
| 301 |
+
.b8 114
|
| 302 |
+
.b8 95
|
| 303 |
+
.b8 102
|
| 304 |
+
.b8 117
|
| 305 |
+
.b8 115
|
| 306 |
+
.b8 101
|
| 307 |
+
.b8 100
|
| 308 |
+
.b8 95
|
| 309 |
+
.b8 99
|
| 310 |
+
.b8 108
|
| 311 |
+
.b8 97
|
| 312 |
+
.b8 109
|
| 313 |
+
.b8 112
|
| 314 |
+
.b8 95
|
| 315 |
+
.b8 109
|
| 316 |
+
.b8 105
|
| 317 |
+
.b8 110
|
| 318 |
+
.b8 95
|
| 319 |
+
.b8 100
|
| 320 |
+
.b8 105
|
| 321 |
+
.b8 118
|
| 322 |
+
.b8 95
|
| 323 |
+
.b8 101
|
| 324 |
+
.b8 113
|
| 325 |
+
.b8 95
|
| 326 |
+
.b8 109
|
| 327 |
+
.b8 117
|
| 328 |
+
.b8 108
|
| 329 |
+
.b8 95
|
| 330 |
+
.b8 115
|
| 331 |
+
.b8 113
|
| 332 |
+
.b8 117
|
| 333 |
+
.b8 101
|
| 334 |
+
.b8 101
|
| 335 |
+
.b8 122
|
| 336 |
+
.b8 101
|
| 337 |
+
.b8 95
|
| 338 |
+
.b8 115
|
| 339 |
+
.b8 117
|
| 340 |
+
.b8 109
|
| 341 |
+
.b8 95
|
| 342 |
+
.b8 52
|
| 343 |
+
.b8 0
|
| 344 |
+
.b8 1 // DW_AT_inline
|
| 345 |
+
.b8 3 // Abbrev [3] 0xc1:0x5e DW_TAG_subprogram
|
| 346 |
+
.b64 $L__func_begin0 // DW_AT_low_pc
|
| 347 |
+
.b64 $L__func_end0 // DW_AT_high_pc
|
| 348 |
+
.b32 139 // DW_AT_abstract_origin
|
| 349 |
+
.b8 4 // Abbrev [4] 0xd6:0x18 DW_TAG_inlined_subroutine
|
| 350 |
+
.b32 139 // DW_AT_abstract_origin
|
| 351 |
+
.b64 $L__tmp1 // DW_AT_low_pc
|
| 352 |
+
.b64 $L__tmp2 // DW_AT_high_pc
|
| 353 |
+
.b8 1 // DW_AT_call_file
|
| 354 |
+
.b8 36 // DW_AT_call_line
|
| 355 |
+
.b8 24 // DW_AT_call_column
|
| 356 |
+
.b8 4 // Abbrev [4] 0xee:0x18 DW_TAG_inlined_subroutine
|
| 357 |
+
.b32 139 // DW_AT_abstract_origin
|
| 358 |
+
.b64 $L__tmp2 // DW_AT_low_pc
|
| 359 |
+
.b64 $L__tmp3 // DW_AT_high_pc
|
| 360 |
+
.b8 1 // DW_AT_call_file
|
| 361 |
+
.b8 38 // DW_AT_call_line
|
| 362 |
+
.b8 24 // DW_AT_call_column
|
| 363 |
+
.b8 4 // Abbrev [4] 0x106:0x18 DW_TAG_inlined_subroutine
|
| 364 |
+
.b32 139 // DW_AT_abstract_origin
|
| 365 |
+
.b64 $L__tmp4 // DW_AT_low_pc
|
| 366 |
+
.b64 $L__tmp5 // DW_AT_high_pc
|
| 367 |
+
.b8 1 // DW_AT_call_file
|
| 368 |
+
.b8 42 // DW_AT_call_line
|
| 369 |
+
.b8 41 // DW_AT_call_column
|
| 370 |
+
.b8 0 // End Of Children Mark
|
| 371 |
+
.b8 0 // End Of Children Mark
|
| 372 |
+
}
|
| 373 |
+
.section .debug_macinfo { }
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.source
ADDED
|
@@ -0,0 +1,206 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":18:0)
|
| 2 |
+
#loc30 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":285:0)
|
| 3 |
+
#loc32 = loc(unknown)
|
| 4 |
+
#loc35 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":260:0)
|
| 5 |
+
#loc39 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":109:0)
|
| 6 |
+
#loc48 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":86:0)
|
| 7 |
+
#loc52 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":63:0)
|
| 8 |
+
#loc61 = loc("in_ptr0"(#loc))
|
| 9 |
+
#loc62 = loc("in_ptr1"(#loc))
|
| 10 |
+
#loc63 = loc("out_ptr2"(#loc))
|
| 11 |
+
#loc64 = loc("r0_numel"(#loc))
|
| 12 |
+
#loc90 = loc("input"(#loc30))
|
| 13 |
+
#loc91 = loc("a"(#loc35))
|
| 14 |
+
#loc92 = loc("b"(#loc35))
|
| 15 |
+
#loc93 = loc("a"(#loc39))
|
| 16 |
+
#loc94 = loc("b"(#loc39))
|
| 17 |
+
#loc98 = loc("x"(#loc48))
|
| 18 |
+
#loc99 = loc("x"(#loc52))
|
| 19 |
+
module {
|
| 20 |
+
tt.func public @triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4(%in_ptr0: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %out_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr2"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
|
| 21 |
+
%xnumel = arith.constant 1 : i32 loc(#loc65)
|
| 22 |
+
%r0_numel_0 = arith.constant 2 : i32 loc(#loc66)
|
| 23 |
+
%xoffset = tt.get_program_id x : i32 loc(#loc67)
|
| 24 |
+
%xoffset_1 = arith.constant 1 : i32 loc(#loc68)
|
| 25 |
+
%xoffset_2 = arith.constant 1 : i32 loc(#loc68)
|
| 26 |
+
%xoffset_3 = arith.muli %xoffset, %xoffset_2 : i32 loc(#loc68)
|
| 27 |
+
%xindex = tt.make_range {end = 1 : i32, start = 0 : i32} : tensor<1xi32> loc(#loc69)
|
| 28 |
+
%xindex_4 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<1xi32> -> tensor<1x1xi32> loc(#loc70)
|
| 29 |
+
%xindex_5 = tt.splat %xoffset_3 : i32 -> tensor<1x1xi32> loc(#loc71)
|
| 30 |
+
%xindex_6 = arith.addi %xindex_5, %xindex_4 : tensor<1x1xi32> loc(#loc71)
|
| 31 |
+
%xmask = arith.constant true loc(#loc72)
|
| 32 |
+
%xmask_7 = arith.constant dense<true> : tensor<1x2xi1> loc(#loc72)
|
| 33 |
+
%r0_index = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32> loc(#loc73)
|
| 34 |
+
%r0_index_8 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<2xi32> -> tensor<1x2xi32> loc(#loc74)
|
| 35 |
+
%r0_offset = arith.constant 0 : i32 loc(#loc75)
|
| 36 |
+
%r0_mask = arith.constant true loc(#loc76)
|
| 37 |
+
%r0_mask_9 = arith.constant dense<true> : tensor<1x2xi1> loc(#loc76)
|
| 38 |
+
%tmp0 = tt.splat %in_ptr0 : !tt.ptr<i64> -> tensor<1x2x!tt.ptr<i64>> loc(#loc77)
|
| 39 |
+
%tmp0_10 = tt.addptr %tmp0, %r0_index_8 : tensor<1x2x!tt.ptr<i64>>, tensor<1x2xi32> loc(#loc77)
|
| 40 |
+
%tmp0_11 = tt.load %tmp0_10 : tensor<1x2x!tt.ptr<i64>> loc(#loc78)
|
| 41 |
+
%tmp4 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1x2x!tt.ptr<i64>> loc(#loc79)
|
| 42 |
+
%tmp4_12 = tt.addptr %tmp4, %r0_index_8 : tensor<1x2x!tt.ptr<i64>>, tensor<1x2xi32> loc(#loc79)
|
| 43 |
+
%tmp4_13 = tt.load %tmp4_12 : tensor<1x2x!tt.ptr<i64>> loc(#loc80)
|
| 44 |
+
%tmp3 = tt.call @"triton.language.standard.sum__i64S1_2S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%tmp0_11) : (tensor<1x2xi64>) -> tensor<1xi64> loc(#loc81)
|
| 45 |
+
%tmp3_14 = tt.expand_dims %tmp3 {axis = 1 : i32} : tensor<1xi64> -> tensor<1x1xi64> loc(#loc82)
|
| 46 |
+
%tmp7 = tt.call @"triton.language.standard.sum__i64S1_2S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%tmp4_13) : (tensor<1x2xi64>) -> tensor<1xi64> loc(#loc83)
|
| 47 |
+
%tmp7_15 = tt.expand_dims %tmp7 {axis = 1 : i32} : tensor<1xi64> -> tensor<1x1xi64> loc(#loc84)
|
| 48 |
+
%tmp8 = arith.sitofp %tmp3_14 : tensor<1x1xi64> to tensor<1x1xf32> loc(#loc85)
|
| 49 |
+
%tmp9 = arith.sitofp %tmp7_15 : tensor<1x1xi64> to tensor<1x1xf32> loc(#loc86)
|
| 50 |
+
%tmp10 = arith.constant 9.99999997E-7 : f32 loc(#loc87)
|
| 51 |
+
%tmp11 = tt.call @torch._inductor.runtime.triton_helpers.maximum__fp32S1_1S_fp32__(%tmp9, %tmp10) : (tensor<1x1xf32>, f32) -> tensor<1x1xf32> loc(#loc88)
|
| 52 |
+
%tmp12 = arith.divf %tmp8, %tmp11 : tensor<1x1xf32> loc(#loc89)
|
| 53 |
+
%c0_i32 = arith.constant 0 : i32 loc(#loc26)
|
| 54 |
+
%cst = arith.constant dense<0> : tensor<1x1xi32> loc(#loc26)
|
| 55 |
+
%0 = tt.splat %out_ptr2 : !tt.ptr<f32> -> tensor<1x1x!tt.ptr<f32>> loc(#loc27)
|
| 56 |
+
%1 = tt.addptr %0, %cst : tensor<1x1x!tt.ptr<f32>>, tensor<1x1xi32> loc(#loc27)
|
| 57 |
+
tt.store %1, %tmp12 : tensor<1x1x!tt.ptr<f32>> loc(#loc28)
|
| 58 |
+
tt.return loc(#loc29)
|
| 59 |
+
} loc(#loc)
|
| 60 |
+
tt.func private @"triton.language.standard.sum__i64S1_2S__(1,)cconstexpr_1__(2,)cconstexpr_False__(3,)cNone"(%input: tensor<1x2xi64> loc("input"(#loc30))) -> tensor<1xi64> attributes {noinline = false} {
|
| 61 |
+
%0 = "tt.reduce"(%input) <{axis = 1 : i32}> ({
|
| 62 |
+
^bb0(%arg1: i64 loc(unknown), %arg2: i64 loc(unknown)):
|
| 63 |
+
%2 = tt.call @triton.language.standard._sum_combine__i64_i64__(%arg1, %arg2) : (i64, i64) -> i64 loc(#loc31)
|
| 64 |
+
tt.reduce.return %2 : i64 loc(#loc31)
|
| 65 |
+
}) : (tensor<1x2xi64>) -> tensor<1xi64> loc(#loc31)
|
| 66 |
+
tt.return %0 : tensor<1xi64> loc(#loc33)
|
| 67 |
+
^bb1: // no predecessors
|
| 68 |
+
%1 = ub.poison : tensor<1xi64> loc(#loc34)
|
| 69 |
+
tt.return %1 : tensor<1xi64> loc(#loc34)
|
| 70 |
+
} loc(#loc30)
|
| 71 |
+
tt.func private @triton.language.standard._sum_combine__i64_i64__(%a: i64 loc("a"(#loc35)), %b: i64 loc("b"(#loc35))) -> i64 attributes {noinline = false} {
|
| 72 |
+
%0 = arith.addi %a, %b : i64 loc(#loc36)
|
| 73 |
+
tt.return %0 : i64 loc(#loc37)
|
| 74 |
+
^bb1: // no predecessors
|
| 75 |
+
%1 = ub.poison : i64 loc(#loc38)
|
| 76 |
+
tt.return %1 : i64 loc(#loc38)
|
| 77 |
+
} loc(#loc35)
|
| 78 |
+
tt.func private @torch._inductor.runtime.triton_helpers.maximum__fp32S1_1S_fp32__(%a: tensor<1x1xf32> loc("a"(#loc39)), %b: f32 loc("b"(#loc39))) -> tensor<1x1xf32> attributes {noinline = false} {
|
| 79 |
+
%mask = tt.splat %b : f32 -> tensor<1x1xf32> loc(#loc95)
|
| 80 |
+
%mask_0 = arith.cmpf ogt, %a, %mask : tensor<1x1xf32> loc(#loc100)
|
| 81 |
+
%0 = tt.call @torch._inductor.runtime.triton_helpers.is_floating__fp32S1_1S__(%a) : (tensor<1x1xf32>) -> i1 loc(#loc41)
|
| 82 |
+
%1 = scf.if %0 -> (tensor<1x1xi1>) {
|
| 83 |
+
%mask_1 = arith.cmpf une, %a, %a : tensor<1x1xf32> loc(#loc96)
|
| 84 |
+
%mask_2 = arith.ori %mask_0, %mask_1 : tensor<1x1xi1> loc(#loc101)
|
| 85 |
+
scf.yield %mask_2 : tensor<1x1xi1> loc(#loc101)
|
| 86 |
+
} else {
|
| 87 |
+
scf.yield %mask_0 : tensor<1x1xi1> loc(#loc32)
|
| 88 |
+
} loc(#loc42)
|
| 89 |
+
%2 = tt.splat %b : f32 -> tensor<1x1xf32> loc(#loc45)
|
| 90 |
+
%3 = arith.select %1, %a, %2 : tensor<1x1xi1>, tensor<1x1xf32> loc(#loc45)
|
| 91 |
+
tt.return %3 : tensor<1x1xf32> loc(#loc46)
|
| 92 |
+
^bb1: // no predecessors
|
| 93 |
+
%4 = ub.poison : tensor<1x1xf32> loc(#loc47)
|
| 94 |
+
tt.return %4 : tensor<1x1xf32> loc(#loc47)
|
| 95 |
+
} loc(#loc39)
|
| 96 |
+
tt.func private @torch._inductor.runtime.triton_helpers.is_floating__fp32S1_1S__(%x: tensor<1x1xf32> loc("x"(#loc48))) -> i1 attributes {noinline = false} {
|
| 97 |
+
%0 = tt.call @torch._inductor.runtime.triton_helpers.promote_to_tensor__fp32S1_1S__(%x) : (tensor<1x1xf32>) -> tensor<1x1xf32> loc(#loc49)
|
| 98 |
+
%true = arith.constant true loc(#loc50)
|
| 99 |
+
tt.return %true : i1 loc(#loc50)
|
| 100 |
+
^bb1: // no predecessors
|
| 101 |
+
%1 = ub.poison : i1 loc(#loc51)
|
| 102 |
+
tt.return %1 : i1 loc(#loc51)
|
| 103 |
+
} loc(#loc48)
|
| 104 |
+
tt.func private @torch._inductor.runtime.triton_helpers.promote_to_tensor__fp32S1_1S__(%x: tensor<1x1xf32> loc("x"(#loc52))) -> tensor<1x1xf32> attributes {noinline = false} {
|
| 105 |
+
%0 = tt.call @"triton.language.standard.zeros____(0, 0)cconstexpr_1__(1,)cconstexpr_int1_"() : () -> tensor<1xi1> loc(#loc53)
|
| 106 |
+
%1 = arith.uitofp %0 : tensor<1xi1> to tensor<1xf32> loc(#loc54)
|
| 107 |
+
%2 = tt.expand_dims %1 {axis = 0 : i32} : tensor<1xf32> -> tensor<1x1xf32> loc(#loc54)
|
| 108 |
+
%3 = arith.addf %x, %2 : tensor<1x1xf32> loc(#loc54)
|
| 109 |
+
tt.return %3 : tensor<1x1xf32> loc(#loc55)
|
| 110 |
+
^bb1: // no predecessors
|
| 111 |
+
%4 = ub.poison : tensor<1x1xf32> loc(#loc56)
|
| 112 |
+
tt.return %4 : tensor<1x1xf32> loc(#loc56)
|
| 113 |
+
} loc(#loc52)
|
| 114 |
+
tt.func private @"triton.language.standard.zeros____(0, 0)cconstexpr_1__(1,)cconstexpr_int1_"() -> tensor<1xi1> attributes {noinline = false} {
|
| 115 |
+
%false = arith.constant false loc(#loc58)
|
| 116 |
+
%cst = arith.constant dense<false> : tensor<1xi1> loc(#loc58)
|
| 117 |
+
tt.return %cst : tensor<1xi1> loc(#loc59)
|
| 118 |
+
^bb1: // no predecessors
|
| 119 |
+
%0 = ub.poison : tensor<1xi1> loc(#loc60)
|
| 120 |
+
tt.return %0 : tensor<1xi1> loc(#loc60)
|
| 121 |
+
} loc(#loc57)
|
| 122 |
+
} loc(#loc)
|
| 123 |
+
#loc1 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":19:13)
|
| 124 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":20:15)
|
| 125 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":24:28)
|
| 126 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":24:33)
|
| 127 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":25:36)
|
| 128 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":25:44)
|
| 129 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":25:23)
|
| 130 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":26:46)
|
| 131 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":27:28)
|
| 132 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":27:38)
|
| 133 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":28:16)
|
| 134 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":29:48)
|
| 135 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":33:30)
|
| 136 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":33:37)
|
| 137 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":34:30)
|
| 138 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":34:37)
|
| 139 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":36:24)
|
| 140 |
+
#loc18 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":36:27)
|
| 141 |
+
#loc19 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":38:24)
|
| 142 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":38:27)
|
| 143 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":39:19)
|
| 144 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":40:19)
|
| 145 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":41:12)
|
| 146 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":42:41)
|
| 147 |
+
#loc25 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":43:20)
|
| 148 |
+
#loc26 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:49)
|
| 149 |
+
#loc27 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:25)
|
| 150 |
+
#loc28 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:68)
|
| 151 |
+
#loc29 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:4)
|
| 152 |
+
#loc31 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 153 |
+
#loc33 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:11)
|
| 154 |
+
#loc34 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:4)
|
| 155 |
+
#loc36 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 156 |
+
#loc37 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:11)
|
| 157 |
+
#loc38 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:4)
|
| 158 |
+
#loc40 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":110:15)
|
| 159 |
+
#loc41 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":111:19)
|
| 160 |
+
#loc42 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":111:7)
|
| 161 |
+
#loc43 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:21)
|
| 162 |
+
#loc44 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:16)
|
| 163 |
+
#loc45 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":113:29)
|
| 164 |
+
#loc46 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":113:11)
|
| 165 |
+
#loc47 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":113:4)
|
| 166 |
+
#loc49 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":87:29)
|
| 167 |
+
#loc50 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":87:11)
|
| 168 |
+
#loc51 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":87:4)
|
| 169 |
+
#loc53 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":65:30)
|
| 170 |
+
#loc54 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":65:15)
|
| 171 |
+
#loc55 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":65:11)
|
| 172 |
+
#loc56 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":65:4)
|
| 173 |
+
#loc57 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":118:0)
|
| 174 |
+
#loc58 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":127:31)
|
| 175 |
+
#loc59 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":127:11)
|
| 176 |
+
#loc60 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":127:4)
|
| 177 |
+
#loc65 = loc("xnumel"(#loc1))
|
| 178 |
+
#loc66 = loc("r0_numel"(#loc2))
|
| 179 |
+
#loc67 = loc("xoffset"(#loc3))
|
| 180 |
+
#loc68 = loc("xoffset"(#loc4))
|
| 181 |
+
#loc69 = loc("xindex"(#loc5))
|
| 182 |
+
#loc70 = loc("xindex"(#loc6))
|
| 183 |
+
#loc71 = loc("xindex"(#loc7))
|
| 184 |
+
#loc72 = loc("xmask"(#loc8))
|
| 185 |
+
#loc73 = loc("r0_index"(#loc9))
|
| 186 |
+
#loc74 = loc("r0_index"(#loc10))
|
| 187 |
+
#loc75 = loc("r0_offset"(#loc11))
|
| 188 |
+
#loc76 = loc("r0_mask"(#loc12))
|
| 189 |
+
#loc77 = loc("tmp0"(#loc13))
|
| 190 |
+
#loc78 = loc("tmp0"(#loc14))
|
| 191 |
+
#loc79 = loc("tmp4"(#loc15))
|
| 192 |
+
#loc80 = loc("tmp4"(#loc16))
|
| 193 |
+
#loc81 = loc("tmp3"(#loc17))
|
| 194 |
+
#loc82 = loc("tmp3"(#loc18))
|
| 195 |
+
#loc83 = loc("tmp7"(#loc19))
|
| 196 |
+
#loc84 = loc("tmp7"(#loc20))
|
| 197 |
+
#loc85 = loc("tmp8"(#loc21))
|
| 198 |
+
#loc86 = loc("tmp9"(#loc22))
|
| 199 |
+
#loc87 = loc("tmp10"(#loc23))
|
| 200 |
+
#loc88 = loc("tmp11"(#loc24))
|
| 201 |
+
#loc89 = loc("tmp12"(#loc25))
|
| 202 |
+
#loc95 = loc("mask"(#loc40))
|
| 203 |
+
#loc96 = loc("mask"(#loc43))
|
| 204 |
+
#loc97 = loc("mask"(#loc44))
|
| 205 |
+
#loc100 = loc("mask"(#loc95))
|
| 206 |
+
#loc101 = loc("mask"(#loc97))
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttgir
ADDED
|
@@ -0,0 +1,91 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#blocked = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [1, 32], warpsPerCTA = [1, 2], order = [0, 1]}>
|
| 2 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":18:0)
|
| 3 |
+
#loc1 = loc(unknown)
|
| 4 |
+
#loc8 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":36:24)
|
| 5 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":38:24)
|
| 6 |
+
#loc23 = loc("in_ptr0"(#loc))
|
| 7 |
+
#loc24 = loc("in_ptr1"(#loc))
|
| 8 |
+
#loc25 = loc("out_ptr2"(#loc))
|
| 9 |
+
#loc26 = loc("r0_numel"(#loc))
|
| 10 |
+
#loc32 = loc("tmp3"(#loc8))
|
| 11 |
+
#loc34 = loc("tmp7"(#loc11))
|
| 12 |
+
#loc44 = loc(callsite(#loc1 at #loc32))
|
| 13 |
+
#loc46 = loc(callsite(#loc1 at #loc34))
|
| 14 |
+
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 2 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
|
| 15 |
+
tt.func public @triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4(%in_ptr0: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %out_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr2"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
|
| 16 |
+
%cst = arith.constant dense<9.99999997E-7> : tensor<1x1xf32, #blocked> loc(#loc1)
|
| 17 |
+
%r0_index = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc27)
|
| 18 |
+
%r0_index_0 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<2xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x2xi32, #blocked> loc(#loc27)
|
| 19 |
+
%tmp0 = tt.splat %in_ptr0 : !tt.ptr<i64> -> tensor<1x2x!tt.ptr<i64>, #blocked> loc(#loc28)
|
| 20 |
+
%tmp0_1 = tt.addptr %tmp0, %r0_index_0 : tensor<1x2x!tt.ptr<i64>, #blocked>, tensor<1x2xi32, #blocked> loc(#loc28)
|
| 21 |
+
%tmp0_2 = tt.load %tmp0_1 : tensor<1x2x!tt.ptr<i64>, #blocked> loc(#loc29)
|
| 22 |
+
%tmp4 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1x2x!tt.ptr<i64>, #blocked> loc(#loc30)
|
| 23 |
+
%tmp4_3 = tt.addptr %tmp4, %r0_index_0 : tensor<1x2x!tt.ptr<i64>, #blocked>, tensor<1x2xi32, #blocked> loc(#loc30)
|
| 24 |
+
%tmp4_4 = tt.load %tmp4_3 : tensor<1x2x!tt.ptr<i64>, #blocked> loc(#loc31)
|
| 25 |
+
%tmp3 = "tt.reduce"(%tmp0_2) <{axis = 1 : i32}> ({
|
| 26 |
+
^bb0(%tmp3_9: i64 loc(callsite(#loc1 at #loc32)), %tmp3_10: i64 loc(callsite(#loc1 at #loc32))):
|
| 27 |
+
%tmp3_11 = arith.addi %tmp3_9, %tmp3_10 : i64 loc(#loc51)
|
| 28 |
+
tt.reduce.return %tmp3_11 : i64 loc(#loc43)
|
| 29 |
+
}) : (tensor<1x2xi64, #blocked>) -> tensor<1xi64, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc43)
|
| 30 |
+
%tmp3_5 = tt.expand_dims %tmp3 {axis = 1 : i32} : tensor<1xi64, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<1x1xi64, #blocked> loc(#loc33)
|
| 31 |
+
%tmp7 = "tt.reduce"(%tmp4_4) <{axis = 1 : i32}> ({
|
| 32 |
+
^bb0(%tmp7_9: i64 loc(callsite(#loc1 at #loc34)), %tmp7_10: i64 loc(callsite(#loc1 at #loc34))):
|
| 33 |
+
%tmp7_11 = arith.addi %tmp7_9, %tmp7_10 : i64 loc(#loc52)
|
| 34 |
+
tt.reduce.return %tmp7_11 : i64 loc(#loc45)
|
| 35 |
+
}) : (tensor<1x2xi64, #blocked>) -> tensor<1xi64, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc45)
|
| 36 |
+
%tmp7_6 = tt.expand_dims %tmp7 {axis = 1 : i32} : tensor<1xi64, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<1x1xi64, #blocked> loc(#loc35)
|
| 37 |
+
%tmp8 = arith.sitofp %tmp3_5 : tensor<1x1xi64, #blocked> to tensor<1x1xf32, #blocked> loc(#loc36)
|
| 38 |
+
%tmp9 = arith.sitofp %tmp7_6 : tensor<1x1xi64, #blocked> to tensor<1x1xf32, #blocked> loc(#loc37)
|
| 39 |
+
%mask = arith.cmpf ogt, %tmp9, %cst : tensor<1x1xf32, #blocked> loc(#loc53)
|
| 40 |
+
%mask_7 = arith.cmpf une, %tmp9, %tmp9 : tensor<1x1xf32, #blocked> loc(#loc48)
|
| 41 |
+
%mask_8 = arith.ori %mask, %mask_7 : tensor<1x1xi1, #blocked> loc(#loc54)
|
| 42 |
+
%tmp11 = arith.select %mask_8, %tmp9, %cst : tensor<1x1xi1, #blocked>, tensor<1x1xf32, #blocked> loc(#loc50)
|
| 43 |
+
%tmp12 = arith.divf %tmp8, %tmp11 : tensor<1x1xf32, #blocked> loc(#loc42)
|
| 44 |
+
%0 = tt.splat %out_ptr2 : !tt.ptr<f32> -> tensor<1x1x!tt.ptr<f32>, #blocked> loc(#loc21)
|
| 45 |
+
tt.store %0, %tmp12 : tensor<1x1x!tt.ptr<f32>, #blocked> loc(#loc21)
|
| 46 |
+
tt.return loc(#loc22)
|
| 47 |
+
} loc(#loc)
|
| 48 |
+
} loc(#loc)
|
| 49 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":27:38)
|
| 50 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":33:30)
|
| 51 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":33:37)
|
| 52 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":34:30)
|
| 53 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":34:37)
|
| 54 |
+
#loc7 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 55 |
+
#loc9 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 56 |
+
#loc10 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":36:27)
|
| 57 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":38:27)
|
| 58 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":39:19)
|
| 59 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":40:19)
|
| 60 |
+
#loc15 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":110:15)
|
| 61 |
+
#loc16 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":42:41)
|
| 62 |
+
#loc17 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:21)
|
| 63 |
+
#loc18 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:16)
|
| 64 |
+
#loc19 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":113:29)
|
| 65 |
+
#loc20 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":43:20)
|
| 66 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:68)
|
| 67 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:4)
|
| 68 |
+
#loc27 = loc("r0_index"(#loc2))
|
| 69 |
+
#loc28 = loc("tmp0"(#loc3))
|
| 70 |
+
#loc29 = loc("tmp0"(#loc4))
|
| 71 |
+
#loc30 = loc("tmp4"(#loc5))
|
| 72 |
+
#loc31 = loc("tmp4"(#loc6))
|
| 73 |
+
#loc33 = loc("tmp3"(#loc10))
|
| 74 |
+
#loc35 = loc("tmp7"(#loc12))
|
| 75 |
+
#loc36 = loc("tmp8"(#loc13))
|
| 76 |
+
#loc37 = loc("tmp9"(#loc14))
|
| 77 |
+
#loc38 = loc("mask"(#loc15))
|
| 78 |
+
#loc39 = loc("tmp11"(#loc16))
|
| 79 |
+
#loc40 = loc("mask"(#loc17))
|
| 80 |
+
#loc41 = loc("mask"(#loc18))
|
| 81 |
+
#loc42 = loc("tmp12"(#loc20))
|
| 82 |
+
#loc43 = loc(callsite(#loc7 at #loc32))
|
| 83 |
+
#loc45 = loc(callsite(#loc7 at #loc34))
|
| 84 |
+
#loc47 = loc("mask"(#loc38))
|
| 85 |
+
#loc48 = loc(callsite(#loc40 at #loc39))
|
| 86 |
+
#loc49 = loc("mask"(#loc41))
|
| 87 |
+
#loc50 = loc(callsite(#loc19 at #loc39))
|
| 88 |
+
#loc51 = loc(callsite(#loc9 at #loc43))
|
| 89 |
+
#loc52 = loc(callsite(#loc9 at #loc45))
|
| 90 |
+
#loc53 = loc(callsite(#loc47 at #loc39))
|
| 91 |
+
#loc54 = loc(callsite(#loc49 at #loc39))
|
SpecForge-ext/cache/compiled_kernels/triton/1/5QCATXXJOHVG6NQ75K46KBBWQRORJ7QMSXCTH7GV3JOSAQPCCD4Q/triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4.ttir
ADDED
|
@@ -0,0 +1,93 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#loc = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":18:0)
|
| 2 |
+
#loc1 = loc(unknown)
|
| 3 |
+
#loc9 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":36:24)
|
| 4 |
+
#loc12 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":38:24)
|
| 5 |
+
#loc25 = loc("in_ptr0"(#loc))
|
| 6 |
+
#loc26 = loc("in_ptr1"(#loc))
|
| 7 |
+
#loc27 = loc("out_ptr2"(#loc))
|
| 8 |
+
#loc28 = loc("r0_numel"(#loc))
|
| 9 |
+
#loc35 = loc("tmp3"(#loc9))
|
| 10 |
+
#loc37 = loc("tmp7"(#loc12))
|
| 11 |
+
#loc47 = loc(callsite(#loc1 at #loc35))
|
| 12 |
+
#loc49 = loc(callsite(#loc1 at #loc37))
|
| 13 |
+
module {
|
| 14 |
+
tt.func public @triton_per_fused_clamp_min_div_eq_mul_squeeze_sum_4(%in_ptr0: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr0"(#loc)), %in_ptr1: !tt.ptr<i64> {tt.divisibility = 16 : i32} loc("in_ptr1"(#loc)), %out_ptr2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("out_ptr2"(#loc)), %r0_numel: i32 loc("r0_numel"(#loc))) attributes {noinline = false} {
|
| 15 |
+
%cst = arith.constant dense<9.99999997E-7> : tensor<1x1xf32> loc(#loc1)
|
| 16 |
+
%r0_index = tt.make_range {end = 2 : i32, start = 0 : i32} : tensor<2xi32> loc(#loc29)
|
| 17 |
+
%r0_index_0 = tt.expand_dims %r0_index {axis = 0 : i32} : tensor<2xi32> -> tensor<1x2xi32> loc(#loc30)
|
| 18 |
+
%tmp0 = tt.splat %in_ptr0 : !tt.ptr<i64> -> tensor<1x2x!tt.ptr<i64>> loc(#loc31)
|
| 19 |
+
%tmp0_1 = tt.addptr %tmp0, %r0_index_0 : tensor<1x2x!tt.ptr<i64>>, tensor<1x2xi32> loc(#loc31)
|
| 20 |
+
%tmp0_2 = tt.load %tmp0_1 : tensor<1x2x!tt.ptr<i64>> loc(#loc32)
|
| 21 |
+
%tmp4 = tt.splat %in_ptr1 : !tt.ptr<i64> -> tensor<1x2x!tt.ptr<i64>> loc(#loc33)
|
| 22 |
+
%tmp4_3 = tt.addptr %tmp4, %r0_index_0 : tensor<1x2x!tt.ptr<i64>>, tensor<1x2xi32> loc(#loc33)
|
| 23 |
+
%tmp4_4 = tt.load %tmp4_3 : tensor<1x2x!tt.ptr<i64>> loc(#loc34)
|
| 24 |
+
%tmp3 = "tt.reduce"(%tmp0_2) <{axis = 1 : i32}> ({
|
| 25 |
+
^bb0(%tmp3_9: i64 loc(callsite(#loc1 at #loc35)), %tmp3_10: i64 loc(callsite(#loc1 at #loc35))):
|
| 26 |
+
%tmp3_11 = arith.addi %tmp3_9, %tmp3_10 : i64 loc(#loc54)
|
| 27 |
+
tt.reduce.return %tmp3_11 : i64 loc(#loc46)
|
| 28 |
+
}) : (tensor<1x2xi64>) -> tensor<1xi64> loc(#loc46)
|
| 29 |
+
%tmp3_5 = tt.expand_dims %tmp3 {axis = 1 : i32} : tensor<1xi64> -> tensor<1x1xi64> loc(#loc36)
|
| 30 |
+
%tmp7 = "tt.reduce"(%tmp4_4) <{axis = 1 : i32}> ({
|
| 31 |
+
^bb0(%tmp7_9: i64 loc(callsite(#loc1 at #loc37)), %tmp7_10: i64 loc(callsite(#loc1 at #loc37))):
|
| 32 |
+
%tmp7_11 = arith.addi %tmp7_9, %tmp7_10 : i64 loc(#loc55)
|
| 33 |
+
tt.reduce.return %tmp7_11 : i64 loc(#loc48)
|
| 34 |
+
}) : (tensor<1x2xi64>) -> tensor<1xi64> loc(#loc48)
|
| 35 |
+
%tmp7_6 = tt.expand_dims %tmp7 {axis = 1 : i32} : tensor<1xi64> -> tensor<1x1xi64> loc(#loc38)
|
| 36 |
+
%tmp8 = arith.sitofp %tmp3_5 : tensor<1x1xi64> to tensor<1x1xf32> loc(#loc39)
|
| 37 |
+
%tmp9 = arith.sitofp %tmp7_6 : tensor<1x1xi64> to tensor<1x1xf32> loc(#loc40)
|
| 38 |
+
%mask = arith.cmpf ogt, %tmp9, %cst : tensor<1x1xf32> loc(#loc56)
|
| 39 |
+
%mask_7 = arith.cmpf une, %tmp9, %tmp9 : tensor<1x1xf32> loc(#loc51)
|
| 40 |
+
%mask_8 = arith.ori %mask, %mask_7 : tensor<1x1xi1> loc(#loc57)
|
| 41 |
+
%tmp11 = arith.select %mask_8, %tmp9, %cst : tensor<1x1xi1>, tensor<1x1xf32> loc(#loc53)
|
| 42 |
+
%tmp12 = arith.divf %tmp8, %tmp11 : tensor<1x1xf32> loc(#loc45)
|
| 43 |
+
%0 = tt.splat %out_ptr2 : !tt.ptr<f32> -> tensor<1x1x!tt.ptr<f32>> loc(#loc22)
|
| 44 |
+
tt.store %0, %tmp12 : tensor<1x1x!tt.ptr<f32>> loc(#loc23)
|
| 45 |
+
tt.return loc(#loc24)
|
| 46 |
+
} loc(#loc)
|
| 47 |
+
} loc(#loc)
|
| 48 |
+
#loc2 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":27:28)
|
| 49 |
+
#loc3 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":27:38)
|
| 50 |
+
#loc4 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":33:30)
|
| 51 |
+
#loc5 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":33:37)
|
| 52 |
+
#loc6 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":34:30)
|
| 53 |
+
#loc7 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":34:37)
|
| 54 |
+
#loc8 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":291:36)
|
| 55 |
+
#loc10 = loc("/workspace/specforge/lib/python3.11/site-packages/triton/language/standard.py":261:15)
|
| 56 |
+
#loc11 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":36:27)
|
| 57 |
+
#loc13 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":38:27)
|
| 58 |
+
#loc14 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":39:19)
|
| 59 |
+
#loc15 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":40:19)
|
| 60 |
+
#loc16 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":110:15)
|
| 61 |
+
#loc17 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":42:41)
|
| 62 |
+
#loc18 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:21)
|
| 63 |
+
#loc19 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":112:16)
|
| 64 |
+
#loc20 = loc("/workspace/specforge/lib/python3.11/site-packages/torch/_inductor/runtime/triton_helpers.py":113:29)
|
| 65 |
+
#loc21 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":43:20)
|
| 66 |
+
#loc22 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:25)
|
| 67 |
+
#loc23 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:68)
|
| 68 |
+
#loc24 = loc("/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/h5/ch5qtsfeyzopygfndrf54hocf6rdg72pdj4epzueagf5vz5zsyb7.py":44:4)
|
| 69 |
+
#loc29 = loc("r0_index"(#loc2))
|
| 70 |
+
#loc30 = loc("r0_index"(#loc3))
|
| 71 |
+
#loc31 = loc("tmp0"(#loc4))
|
| 72 |
+
#loc32 = loc("tmp0"(#loc5))
|
| 73 |
+
#loc33 = loc("tmp4"(#loc6))
|
| 74 |
+
#loc34 = loc("tmp4"(#loc7))
|
| 75 |
+
#loc36 = loc("tmp3"(#loc11))
|
| 76 |
+
#loc38 = loc("tmp7"(#loc13))
|
| 77 |
+
#loc39 = loc("tmp8"(#loc14))
|
| 78 |
+
#loc40 = loc("tmp9"(#loc15))
|
| 79 |
+
#loc41 = loc("mask"(#loc16))
|
| 80 |
+
#loc42 = loc("tmp11"(#loc17))
|
| 81 |
+
#loc43 = loc("mask"(#loc18))
|
| 82 |
+
#loc44 = loc("mask"(#loc19))
|
| 83 |
+
#loc45 = loc("tmp12"(#loc21))
|
| 84 |
+
#loc46 = loc(callsite(#loc8 at #loc35))
|
| 85 |
+
#loc48 = loc(callsite(#loc8 at #loc37))
|
| 86 |
+
#loc50 = loc("mask"(#loc41))
|
| 87 |
+
#loc51 = loc(callsite(#loc43 at #loc42))
|
| 88 |
+
#loc52 = loc("mask"(#loc44))
|
| 89 |
+
#loc53 = loc(callsite(#loc20 at #loc42))
|
| 90 |
+
#loc54 = loc(callsite(#loc10 at #loc46))
|
| 91 |
+
#loc55 = loc(callsite(#loc10 at #loc48))
|
| 92 |
+
#loc56 = loc(callsite(#loc50 at #loc42))
|
| 93 |
+
#loc57 = loc(callsite(#loc52 at #loc42))
|
SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/__grp__triton_poi_fused_new_zeros_1.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"child_paths": {"triton_poi_fused_new_zeros_1.source": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.source", "triton_poi_fused_new_zeros_1.ttir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.ttir", "triton_poi_fused_new_zeros_1.ttgir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.ttgir", "triton_poi_fused_new_zeros_1.llir": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.llir", "triton_poi_fused_new_zeros_1.ptx": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.ptx", "triton_poi_fused_new_zeros_1.cubin": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.cubin", "triton_poi_fused_new_zeros_1.json": "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.json"}}
|
SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.cubin
ADDED
|
Binary file (5.38 kB). View file
|
|
|
SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.json
ADDED
|
@@ -0,0 +1 @@
|
|
|
|
|
|
|
| 1 |
+
{"hash": "f143f479009090646c7252c2be4342b94b67e2d7b7c0c3b6941ade655041dbc6", "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_new_zeros_1"}
|
SpecForge-ext/cache/compiled_kernels/triton/1/6FB7I6IASCIGI3DSKLBL4Q2CXFFWPYWXW7AMHNUUDLPGKUCB3PDA/triton_poi_fused_new_zeros_1.llir
ADDED
|
@@ -0,0 +1,47 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
; ModuleID = 'LLVMDialectModule'
|
| 2 |
+
source_filename = "LLVMDialectModule"
|
| 3 |
+
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"
|
| 4 |
+
|
| 5 |
+
; Function Attrs: nounwind
|
| 6 |
+
define ptx_kernel void @triton_poi_fused_new_zeros_1(ptr addrspace(1) %0, i32 %1, ptr addrspace(1) readnone captures(none) %2, ptr addrspace(1) readnone captures(none) %3) local_unnamed_addr #0 !dbg !4 {
|
| 7 |
+
%5 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !7
|
| 8 |
+
%6 = shl i32 %5, 8, !dbg !8
|
| 9 |
+
%7 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !9
|
| 10 |
+
%8 = shl nuw nsw i32 %7, 1, !dbg !9
|
| 11 |
+
%9 = and i32 %8, 254, !dbg !9
|
| 12 |
+
%10 = or disjoint i32 %9, %6, !dbg !10
|
| 13 |
+
%11 = icmp slt i32 %10, 544, !dbg !11
|
| 14 |
+
%12 = sext i32 %10 to i64, !dbg !12
|
| 15 |
+
%13 = getelementptr i32, ptr addrspace(1) %0, i64 %12, !dbg !12
|
| 16 |
+
tail call void asm sideeffect "@$3 st.global.v2.b32 [ $2 + 0 ], { $0, $1 };", "r,r,l,b"(i32 0, i32 0, ptr addrspace(1) %13, i1 %11) #2, !dbg !13
|
| 17 |
+
ret void, !dbg !14
|
| 18 |
+
}
|
| 19 |
+
|
| 20 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 21 |
+
declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1
|
| 22 |
+
|
| 23 |
+
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
|
| 24 |
+
declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
|
| 25 |
+
|
| 26 |
+
attributes #0 = { nounwind "nvvm.reqntid"="128" }
|
| 27 |
+
attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
|
| 28 |
+
attributes #2 = { nounwind }
|
| 29 |
+
|
| 30 |
+
!llvm.dbg.cu = !{!0}
|
| 31 |
+
!llvm.module.flags = !{!2, !3}
|
| 32 |
+
|
| 33 |
+
!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
|
| 34 |
+
!1 = !DIFile(filename: "clqbbuegamppef5fckuka5vp5vs5u35wgqa2yxhsbyvq2waia5y4.py", directory: "/workspace/hanrui/SpecForge-ext/cache/compiled_kernels/lq")
|
| 35 |
+
!2 = !{i32 2, !"Debug Info Version", i32 3}
|
| 36 |
+
!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
|
| 37 |
+
!4 = distinct !DISubprogram(name: "triton_poi_fused_new_zeros_1", linkageName: "triton_poi_fused_new_zeros_1", scope: !1, file: !1, line: 18, type: !5, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
|
| 38 |
+
!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
|
| 39 |
+
!6 = !{}
|
| 40 |
+
!7 = !DILocation(line: 20, column: 28, scope: !4)
|
| 41 |
+
!8 = !DILocation(line: 20, column: 33, scope: !4)
|
| 42 |
+
!9 = !DILocation(line: 21, column: 36, scope: !4)
|
| 43 |
+
!10 = !DILocation(line: 21, column: 23, scope: !4)
|
| 44 |
+
!11 = !DILocation(line: 22, column: 21, scope: !4)
|
| 45 |
+
!12 = !DILocation(line: 25, column: 25, scope: !4)
|
| 46 |
+
!13 = !DILocation(line: 25, column: 36, scope: !4)
|
| 47 |
+
!14 = !DILocation(line: 25, column: 4, scope: !4)
|